From 1531d85f185a0b9d9753e09fbf44993d8dc4a3cd Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 11 Aug 2026 00:34:35 +0000 Subject: [PATCH 01/12] first draft --- cpp/benchmarks/CMakeLists.txt | 5 + .../parquet/experimental/variant/extract.cpp | 247 ++++++++++++++++++ 2 files changed, 252 insertions(+) create mode 100644 cpp/benchmarks/io/parquet/experimental/variant/extract.cpp diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 1b57e3b23666..6b35450cc899 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -351,6 +351,11 @@ target_compile_definitions( ) target_link_libraries(PARQUET_DELETION_VECTORS_NVBENCH PRIVATE roaring) +# ################################################################################################## +# * parquet variant extract benchmark +# ---------------------------------------------------------------------- +ConfigureNVBench(VARIANT_NVBENCH io/parquet/experimental/variant/extract.cpp) + # ################################################################################################## # * parquet multithread reader benchmark # ---------------------------------------------------------------------- diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp new file mode 100644 index 000000000000..348aad6a9e61 --- /dev/null +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -0,0 +1,247 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include +#include + +namespace { + +void append_le(std::vector& out, uint64_t bits, int width) +{ + for (int i = 0; i < width; ++i) { + out.push_back(static_cast((bits >> (8 * i)) & 0xff)); + } +} + +// Build a V1 VARIANT metadata blob for a sorted key dictionary (1-byte offsets). +std::vector build_metadata(std::vector const& keys) +{ + std::vector out{0x01, static_cast(keys.size())}; + uint8_t running = 0; + std::vector offs{0x00}; + for (auto const& k : keys) { + running = static_cast(running + static_cast(k.size())); + offs.push_back(running); + } + out.insert(out.end(), offs.begin(), offs.end()); + for (auto const& k : keys) { + out.insert(out.end(), k.begin(), k.end()); + } + return out; +} + +// Wrap `inner` as the sole field (field id `fid`) of a 1-field VARIANT object. +// Uses 1-byte field_id_size and 1-byte field_offset_size (value_header=0 → header=0x02). +std::vector wrap_in_object(uint8_t fid, std::vector const& inner) +{ + // Format: object_header(1) + num_fields(1) + fid(1) + offset[0]=0(1) + offset[1]=size(1) + data + std::vector out{0x02, 0x01, fid, 0x00, static_cast(inner.size())}; + out.insert(out.end(), inner.begin(), inner.end()); + return out; +} + +// Build the leaf VARIANT value blob for the requested type. +// +// Header byte composition: (physical_type_id << 2) | basic_type +// PRIMITIVE basic_type = 0, so header = physical_type_id << 2 +// SHORT_STRING basic_type = 1, so header = (length << 2) | 1 +// ARRAY basic_type = 3, so header = (value_header << 2) | 3 +// +// Physical type IDs used: +// INT32 = 5 → header 0x14 +// FLOAT32 = 14 → header 0x38 +// BOOL_TRUE= 1 → header 0x04 +std::vector build_leaf_value(std::string const& type_str) +{ + if (type_str == "int32_t") { + std::vector out{0x14}; + append_le(out, 42u, 4); + return out; + } + if (type_str == "float") { + std::vector out{0x38}; + float const f = 1.0f; + uint32_t u; + std::memcpy(&u, &f, 4); + append_le(out, u, 4); + return out; + } + if (type_str == "bool") { + return {0x04}; // BOOLEAN_TRUE + } + if (type_str == "string") { + // Short string "hello" (5 bytes): (5 << 2) | 1 = 0x15 + return {0x15, 'h', 'e', 'l', 'l', 'o'}; + } + // "array": VARIANT array of two INT32 values [42, 99]; element [1] is accessed in the benchmark. + // Array header 0x03: basic_type=ARRAY(3), value_header=0 (1-byte count, 1-byte offsets). + // 2 elements, offsets [0, 5, 10], then INT32(42) and INT32(99) (5 bytes each). + std::vector out{0x03, 0x02, 0x00, 0x05, 0x0a}; + out.push_back(0x14); + append_le(out, 42u, 4); + out.push_back(0x14); + append_le(out, 99u, 4); + return out; +} + +// Build the full hit-row value blob by wrapping the leaf in `nesting` object levels. +// Keys a,b,c,d,e map to field IDs 0,1,2,3,4 in the shared dictionary. +// For path a.b.c.d.e the outermost object uses fid=0 ("a"). +std::vector build_hit_value(std::string const& type_str, int nesting) +{ + auto val = build_leaf_value(type_str); + for (int i = nesting - 1; i >= 0; --i) { + val = wrap_in_object(static_cast(i), val); + } + return val; +} + +// Build a VARIANT struct column (STRUCT, list>) from per-row byte vectors. +std::unique_ptr build_variant_column( + std::vector> const& meta_rows, + std::vector> const& val_rows, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto const n = static_cast(meta_rows.size()); + + auto build_list_col = + [&](std::vector> const& rows) -> std::unique_ptr { + std::vector offsets(n + 1, 0); + std::vector flat; + for (cudf::size_type i = 0; i < n; ++i) { + flat.insert(flat.end(), rows[i].begin(), rows[i].end()); + offsets[i + 1] = static_cast(flat.size()); + } + + auto d_offsets = + rmm::device_buffer{offsets.data(), offsets.size() * sizeof(int32_t), stream, mr}; + auto d_data = rmm::device_buffer{flat.data(), flat.size() * sizeof(uint8_t), stream, mr}; + + auto off_col = std::make_unique( + cudf::data_type{cudf::type_id::INT32}, n + 1, std::move(d_offsets), rmm::device_buffer{}, 0); + auto data_col = std::make_unique(cudf::data_type{cudf::type_id::UINT8}, + static_cast(flat.size()), + std::move(d_data), + rmm::device_buffer{}, + 0); + + return cudf::make_lists_column(n, std::move(off_col), std::move(data_col), 0, {}, stream, mr); + }; + + std::vector> children; + children.emplace_back(build_list_col(meta_rows)); + children.emplace_back(build_list_col(val_rows)); + return cudf::make_structs_column(n, std::move(children), 0, {}, stream, mr); +} + +// Keys for the shared metadata dictionary: a=0, b=1, c=2, d=3, e=4 (already lexicographically +// sorted). +std::vector get_dict_keys(int nesting) +{ + std::vector keys; + keys.reserve(nesting); + for (int i = 0; i < nesting; ++i) { + keys.emplace_back(1, static_cast('a' + i)); + } + return keys; +} + +// Build the JSONPath-like extraction path. +// For nesting=2, type=array: "a.b[1]" +// For nesting=3, type=string: "a.b.c" +// For nesting=0, type=array: "[1]" +std::string get_path(int nesting, bool is_array) +{ + std::string path; + for (int i = 0; i < nesting; ++i) { + if (i > 0) path += '.'; + path += static_cast('a' + i); + } + if (is_array) path += "[1]"; + return path; +} + +cudf::data_type get_target_type(std::string const& type_str) +{ + if (type_str == "float") return cudf::data_type{cudf::type_id::FLOAT32}; + if (type_str == "bool") return cudf::data_type{cudf::type_id::BOOL8}; + if (type_str == "string") return cudf::data_type{cudf::type_id::STRING}; + // "int32_t" and "array" (element access yields INT32) + return cudf::data_type{cudf::type_id::INT32}; +} + +} // namespace + +static void bench_variant_extract(nvbench::state& state) +{ + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const num_rows = static_cast(state.get_int64("num_rows")); + auto const type_str = state.get_string("type"); + auto const nesting = static_cast(state.get_int64("nesting")); + auto const hit_rate = static_cast(state.get_int64("hit_rate")); + + bool const is_array = (type_str == "array"); + + // Build per-row blobs. + // hit_rate% of rows contain the correctly typed value at the target path. + // Miss rows use a VARIANT null (0x00) which resolves to null on any cast or path traversal. + auto const keys = get_dict_keys(nesting); + auto const meta_blob = build_metadata(keys); + auto const hit_val = build_hit_value(type_str, nesting); + // VARIANT null: header 0x00 (physical_type=NULLVAL, basic=PRIMITIVE) + std::vector const miss_val{0x00}; + + std::vector> meta_rows(num_rows, meta_blob); + std::vector> val_rows(num_rows); + for (cudf::size_type i = 0; i < num_rows; ++i) { + val_rows[i] = (static_cast(i % 100) < hit_rate) ? hit_val : miss_val; + } + + auto col = build_variant_column(meta_rows, val_rows, stream, mr); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + + auto const target_type = get_target_type(type_str); + + // For nesting=0 with a non-array type, the variant value IS the leaf primitive; use cast_variant. + // For arrays at any nesting level, or any nesting >= 1, use extract_variant_field with a path. + bool const use_cast_variant = (nesting == 0 && !is_array); + auto const path = use_cast_variant ? std::string{} : get_path(nesting, is_array); + + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { + if (use_cast_variant) { + std::ignore = cudf::io::parquet::experimental::cast_variant( + col->view().child(1), target_type, stream, mr); + } else { + std::ignore = cudf::io::parquet::experimental::extract_variant_field( + col->view(), path, target_type, stream, mr); + } + }); +} + +NVBENCH_BENCH(bench_variant_extract) + .set_name("bench_variant_extract") + .add_int64_axis("num_rows", {32768, 262144, 2097152}) + .add_string_axis("type", {"string", "float", "bool", "int32_t", "array"}) + .add_int64_axis("nesting", {0, 1, 5}) + .add_int64_axis("hit_rate", {20, 80}); From 16a1e45f7dcf6d1c821953388b365f231ae54e8d Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 11 Aug 2026 21:14:49 +0000 Subject: [PATCH 02/12] comments --- .../parquet/experimental/variant/extract.cpp | 36 +++++++++++++++---- 1 file changed, 29 insertions(+), 7 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 348aad6a9e61..e63b962b6d82 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -19,6 +19,7 @@ #include #include #include +#include #include namespace { @@ -113,6 +114,26 @@ std::vector build_hit_value(std::string const& type_str, int nesting) return val; } +// Build the miss-row value blob: a valid VARIANT that won't match the target path or type. +// For extract_variant_field rows: a 1-level object keyed on "z" (field ID = nesting in the +// dictionary), so traversal fails at the first key lookup while the row remains non-null. +// For cast_variant rows (nesting=0, non-array): a different primitive type so the cast returns +// null. +std::vector build_miss_value(int nesting, bool is_array, std::string const& type_str) +{ + if (nesting == 0 && !is_array) { + // Wrong-type primitive for the cast path. + if (type_str == "bool") { + std::vector out{0x14}; + append_le(out, 0u, 4); + return out; + } + return {0x04}; // BOOLEAN_TRUE + } + // "z" is always the last key in the dictionary, at field ID = nesting. + return wrap_in_object(static_cast(nesting), build_leaf_value(type_str)); +} + // Build a VARIANT struct column (STRUCT, list>) from per-row byte vectors. std::unique_ptr build_variant_column( std::vector> const& meta_rows, @@ -143,7 +164,7 @@ std::unique_ptr build_variant_column( rmm::device_buffer{}, 0); - return cudf::make_lists_column(n, std::move(off_col), std::move(data_col), 0, {}, stream, mr); + return cudf::make_lists_column(n, std::move(off_col), std::move(data_col), 0, {}); }; std::vector> children; @@ -152,15 +173,16 @@ std::unique_ptr build_variant_column( return cudf::make_structs_column(n, std::move(children), 0, {}, stream, mr); } -// Keys for the shared metadata dictionary: a=0, b=1, c=2, d=3, e=4 (already lexicographically -// sorted). +// Keys for the shared metadata dictionary: a=0, b=1, ... plus "z" for miss rows. +// "z" is appended last; lexicographic order is preserved. std::vector get_dict_keys(int nesting) { std::vector keys; - keys.reserve(nesting); + keys.reserve(nesting + 1); for (int i = 0; i < nesting; ++i) { keys.emplace_back(1, static_cast('a' + i)); } + keys.emplace_back("z"); return keys; } @@ -204,12 +226,12 @@ static void bench_variant_extract(nvbench::state& state) // Build per-row blobs. // hit_rate% of rows contain the correctly typed value at the target path. - // Miss rows use a VARIANT null (0x00) which resolves to null on any cast or path traversal. + // Miss rows hold a valid VARIANT with a wrong key ("z") or wrong primitive type so extraction + // returns null without short-circuiting on a trivial null input. auto const keys = get_dict_keys(nesting); auto const meta_blob = build_metadata(keys); auto const hit_val = build_hit_value(type_str, nesting); - // VARIANT null: header 0x00 (physical_type=NULLVAL, basic=PRIMITIVE) - std::vector const miss_val{0x00}; + auto const miss_val = build_miss_value(nesting, is_array, type_str); std::vector> meta_rows(num_rows, meta_blob); std::vector> val_rows(num_rows); From d7601886977b97415c88c61fa5bfeed50ef91dad Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 12 Aug 2026 21:20:47 +0000 Subject: [PATCH 03/12] reviews --- .../parquet/experimental/variant/extract.cpp | 172 +++++++++++++++--- 1 file changed, 147 insertions(+), 25 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index e63b962b6d82..454ca05ea8e9 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -18,6 +18,7 @@ #include #include +#include #include #include #include @@ -186,6 +187,46 @@ std::vector get_dict_keys(int nesting) return keys; } +// Keys for the field-count benchmark: "f00", "f01", ..., "f{N-1}" plus "z" for miss rows. +// All sort before "z", maintaining the required lexicographic order. +std::vector get_dict_keys_for_fields(int num_fields) +{ + std::vector keys; + keys.reserve(num_fields + 1); + for (int i = 0; i < num_fields; ++i) { + keys.emplace_back("f" + std::string(i < 10 ? "0" : "") + std::to_string(i)); + } + keys.emplace_back("z"); + return keys; +} + +// Build a flat object with `num_fields` fields using 1-byte field IDs and 1-byte offsets. +// Field `target_fid` holds `inner`; all other fields hold a dummy BOOLEAN_TRUE (0x04). +std::vector build_flat_object(int num_fields, + int target_fid, + std::vector const& inner) +{ + // object_header(1) + num_fields(1) + field_ids(num_fields) + offsets(num_fields+1) + data + std::vector out{0x02, static_cast(num_fields)}; + for (int i = 0; i < num_fields; ++i) { + out.push_back(static_cast(i)); + } + uint8_t running = 0; + for (int i = 0; i < num_fields; ++i) { + out.push_back(running); + running += static_cast(i == target_fid ? inner.size() : 1u); + } + out.push_back(running); // sentinel offset after last field + for (int i = 0; i < num_fields; ++i) { + if (i == target_fid) { + out.insert(out.end(), inner.begin(), inner.end()); + } else { + out.push_back(0x04); // BOOLEAN_TRUE dummy + } + } + return out; +} + // Build the JSONPath-like extraction path. // For nesting=2, type=array: "a.b[1]" // For nesting=3, type=string: "a.b.c" @@ -212,7 +253,58 @@ cudf::data_type get_target_type(std::string const& type_str) } // namespace -static void bench_variant_extract(nvbench::state& state) +// Assign each row randomly as a hit or miss rather than using contiguous strided ranges, +// so the memory access pattern doesn't accidentally favour cache locality. +void fill_val_rows(std::vector>& val_rows, + std::vector const& hit_val, + std::vector const& miss_val, + int hit_rate) +{ + std::mt19937 rng{42}; + std::uniform_int_distribution dist{0, 99}; + for (auto& row : val_rows) { + row = (dist(rng) < hit_rate) ? hit_val : miss_val; + } +} + +// Benchmarks cast_variant: each row's value IS the leaf primitive (no path traversal). +static void bench_variant_cast(nvbench::state& state) +{ + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const num_rows = static_cast(state.get_int64("num_rows")); + auto const type_str = state.get_string("type"); + auto const hit_rate = static_cast(state.get_int64("hit_rate")); + + auto const meta_blob = build_metadata(get_dict_keys(0)); + auto const hit_val = build_leaf_value(type_str); + auto const miss_val = build_miss_value(0, /*is_array=*/false, type_str); + + std::vector> meta_rows(num_rows, meta_blob); + std::vector> val_rows(num_rows); + fill_val_rows(val_rows, hit_val, miss_val, hit_rate); + + auto col = build_variant_column(meta_rows, val_rows, stream, mr); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + + auto const target_type = get_target_type(type_str); + + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { + std::ignore = + cudf::io::parquet::experimental::cast_variant(col->view().child(1), target_type, stream, mr); + }); +} + +NVBENCH_BENCH(bench_variant_cast) + .set_name("bench_variant_cast") + .add_int64_axis("num_rows", {32768, 262144, 2097152}) + .add_string_axis("type", {"string", "float", "bool", "int32_t"}) + .add_int64_axis("hit_rate", {20, 80}); + +// Benchmarks extract_variant_field with varying path depth (nesting >= 1). +static void bench_variant_extract_nesting(nvbench::state& state) { auto stream = cudf::get_default_stream(); auto mr = cudf::get_current_device_resource_ref(); @@ -224,46 +316,76 @@ static void bench_variant_extract(nvbench::state& state) bool const is_array = (type_str == "array"); - // Build per-row blobs. - // hit_rate% of rows contain the correctly typed value at the target path. - // Miss rows hold a valid VARIANT with a wrong key ("z") or wrong primitive type so extraction - // returns null without short-circuiting on a trivial null input. - auto const keys = get_dict_keys(nesting); - auto const meta_blob = build_metadata(keys); + auto const meta_blob = build_metadata(get_dict_keys(nesting)); auto const hit_val = build_hit_value(type_str, nesting); auto const miss_val = build_miss_value(nesting, is_array, type_str); std::vector> meta_rows(num_rows, meta_blob); std::vector> val_rows(num_rows); - for (cudf::size_type i = 0; i < num_rows; ++i) { - val_rows[i] = (static_cast(i % 100) < hit_rate) ? hit_val : miss_val; - } + fill_val_rows(val_rows, hit_val, miss_val, hit_rate); auto col = build_variant_column(meta_rows, val_rows, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); auto const target_type = get_target_type(type_str); - - // For nesting=0 with a non-array type, the variant value IS the leaf primitive; use cast_variant. - // For arrays at any nesting level, or any nesting >= 1, use extract_variant_field with a path. - bool const use_cast_variant = (nesting == 0 && !is_array); - auto const path = use_cast_variant ? std::string{} : get_path(nesting, is_array); + auto const path = get_path(nesting, is_array); state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { - if (use_cast_variant) { - std::ignore = cudf::io::parquet::experimental::cast_variant( - col->view().child(1), target_type, stream, mr); - } else { - std::ignore = cudf::io::parquet::experimental::extract_variant_field( - col->view(), path, target_type, stream, mr); - } + std::ignore = cudf::io::parquet::experimental::extract_variant_field( + col->view(), path, target_type, stream, mr); }); } -NVBENCH_BENCH(bench_variant_extract) - .set_name("bench_variant_extract") +NVBENCH_BENCH(bench_variant_extract_nesting) + .set_name("bench_variant_extract_nesting") .add_int64_axis("num_rows", {32768, 262144, 2097152}) .add_string_axis("type", {"string", "float", "bool", "int32_t", "array"}) - .add_int64_axis("nesting", {0, 1, 5}) + .add_int64_axis("nesting", {1, 5}) + .add_int64_axis("hit_rate", {20, 80}); + +// Benchmarks extract_variant_field on a flat object, varying the total number of fields +// and whether the target field is first or last (probes binary search cost). +// Type is fixed to int32_t to isolate field-lookup overhead. +static void bench_variant_extract_fields(nvbench::state& state) +{ + auto stream = cudf::get_default_stream(); + auto mr = cudf::get_current_device_resource_ref(); + + auto const num_rows = static_cast(state.get_int64("num_rows")); + auto const num_fields = static_cast(state.get_int64("num_fields")); + auto const field_pos_str = state.get_string("field_position"); + auto const hit_rate = static_cast(state.get_int64("hit_rate")); + + int const target_fid = (field_pos_str == "last") ? (num_fields - 1) : 0; + + auto const meta_blob = build_metadata(get_dict_keys_for_fields(num_fields)); + auto const leaf = build_leaf_value("int32_t"); + auto const hit_val = build_flat_object(num_fields, target_fid, leaf); + // Miss: object keyed on "z" (field ID = num_fields), so the lookup fails. + auto const miss_val = wrap_in_object(static_cast(num_fields), leaf); + + std::vector> meta_rows(num_rows, meta_blob); + std::vector> val_rows(num_rows); + fill_val_rows(val_rows, hit_val, miss_val, hit_rate); + + auto col = build_variant_column(meta_rows, val_rows, stream, mr); + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); + + std::string const path = + "f" + std::string(target_fid < 10 ? "0" : "") + std::to_string(target_fid); + auto const target_type = cudf::data_type{cudf::type_id::INT32}; + + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { + std::ignore = cudf::io::parquet::experimental::extract_variant_field( + col->view(), path, target_type, stream, mr); + }); +} + +NVBENCH_BENCH(bench_variant_extract_fields) + .set_name("bench_variant_extract_fields") + .add_int64_axis("num_rows", {32768, 262144, 2097152}) + .add_int64_axis("num_fields", {1, 10, 100}) + .add_string_axis("field_position", {"first", "last"}) .add_int64_axis("hit_rate", {20, 80}); From 82a2d75f3aab3a69c5b2c853ef3f04e84cc6b45b Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 12 Aug 2026 21:43:25 +0000 Subject: [PATCH 04/12] reviews --- .../parquet/experimental/variant/extract.cpp | 28 ++++++-- .../io/experimental/variant_extract_test.cpp | 65 +++++++++++++++++-- 2 files changed, 79 insertions(+), 14 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 454ca05ea8e9..a13e83cd2969 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -32,17 +32,31 @@ void append_le(std::vector& out, uint64_t bits, int width) } } -// Build a V1 VARIANT metadata blob for a sorted key dictionary (1-byte offsets). +// Build a V1 VARIANT metadata blob for a sorted key dictionary. +// Uses 2-byte offsets when the total string length exceeds 255 bytes; 1-byte otherwise. +// Header bits [7:6] = offset_size_minus_one; bits [3:0] = version (1). std::vector build_metadata(std::vector const& keys) { - std::vector out{0x01, static_cast(keys.size())}; - uint8_t running = 0; - std::vector offs{0x00}; + uint32_t total = 0; + for (auto const& k : keys) + total += static_cast(k.size()); + + int const offset_size = (total > 255u) ? 2 : 1; + std::vector out{static_cast(0x01 | ((offset_size - 1) << 6))}; + + auto write_le = [&](uint32_t v) { + for (int i = 0; i < offset_size; ++i) + out.push_back(static_cast(v >> (8 * i))); + }; + write_le(static_cast(keys.size())); + + uint32_t running = 0; + write_le(0u); for (auto const& k : keys) { - running = static_cast(running + static_cast(k.size())); - offs.push_back(running); + running += static_cast(k.size()); + write_le(running); } - out.insert(out.end(), offs.begin(), offs.end()); + for (auto const& k : keys) { out.insert(out.end(), k.begin(), k.end()); } diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 03c334f250e4..f2f79a7e4dfa 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -535,18 +535,30 @@ inline cudf::test::structs_column_wrapper wrap_multi_row_variant( return cudf::test::structs_column_wrapper{std::move(children)}; } -// Build a metadata blob (version 1, offset_size=1) for the given ordered string dictionary. +// Build a V1 VARIANT metadata blob for the given ordered string dictionary. +// Uses 2-byte offsets when total string length exceeds 255 bytes; 1-byte otherwise. +// Header bits [7:6] = offset_size_minus_one; bits [3:0] = version (1). inline std::vector build_metadata(std::vector const& keys) { - std::vector out{0x01, static_cast(keys.size())}; + uint32_t total = 0; + for (auto const& k : keys) + total += static_cast(k.size()); - std::vector offs{0x00}; - uint8_t running = 0; + int const offset_size = (total > 255u) ? 2 : 1; + std::vector out{static_cast(0x01 | ((offset_size - 1) << 6))}; + + auto write_le = [&](uint32_t v) { + for (int i = 0; i < offset_size; ++i) + out.push_back(static_cast(v >> (8 * i))); + }; + write_le(static_cast(keys.size())); + + uint32_t running = 0; + write_le(0u); for (auto const& k : keys) { - running = static_cast(running + k.size()); - offs.push_back(running); + running += static_cast(k.size()); + write_le(running); } - out.insert(out.end(), offs.begin(), offs.end()); for (auto const& k : keys) { out.insert(out.end(), k.begin(), k.end()); @@ -768,6 +780,45 @@ TEST_F(ExtractVariantFieldTest, LargeDictionaryAndObjectScan) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*last, cudf::test::fixed_width_column_wrapper{49}); } +TEST_F(ExtractVariantFieldTest, LargeDictionary100FieldsExtractLast) +{ + // 100-key dictionary "k00"..."k99" totals 300 string bytes (> 255), so build_metadata must emit + // 2-byte offsets. The value is a flat 100-field object where field 99 ("k99") holds INT32(99) + // and all other fields hold BOOLEAN_TRUE (1 byte), keeping value offsets within 1-byte range + // (99 * 1 + 5 = 104 bytes). + auto const keys = make_numeric_keys(100); + auto const meta = build_metadata(keys); + + constexpr int n_fields = 100; + constexpr int target_fid = 99; + auto const target_val = enc_int32(target_fid); + constexpr uint8_t bool_true_byte = 0x04; + + std::vector val{make_variant_object_header(), static_cast(n_fields)}; + for (int i = 0; i < n_fields; ++i) + val.push_back(static_cast(i)); // field IDs + uint8_t off = 0; + for (int i = 0; i < n_fields; ++i) { + val.push_back(off); + off = static_cast(off + (i == target_fid ? target_val.size() : 1u)); + } + val.push_back(off); // sentinel offset after last field + for (int i = 0; i < n_fields; ++i) { + if (i == target_fid) { + val.insert(val.end(), target_val.begin(), target_val.end()); + } else { + val.push_back(bool_true_byte); + } + } + + auto col = wrap_single_variant(meta, val); + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, "k99", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + + cudf::test::fixed_width_column_wrapper expected{int32_t{99}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + TEST_F(ExtractVariantFieldTest, MalformedVariantDataYieldsNull) { // The column shape is a valid STRUCT, list>, but the VARIANT bytes are From 81889d1724e1bc14c054341577c75eac4d4751d0 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 12 Aug 2026 22:15:53 +0000 Subject: [PATCH 05/12] addressing comments --- .../io/experimental/variant_extract_test.cpp | 58 +++++++++++++++++++ 1 file changed, 58 insertions(+) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index f2f79a7e4dfa..69ae171ccaaa 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -819,6 +819,64 @@ TEST_F(ExtractVariantFieldTest, LargeDictionary100FieldsExtractLast) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } +TEST_F(ExtractVariantFieldTest, MetadataOffsetSizeThresholdBoundary) +{ + // Verifies build_metadata selects 1-byte offsets when total string bytes == 255 (still fits) + // and 2-byte offsets when total == 256 (first value that overflows a uint8_t accumulator). + auto stream = cudf::test::get_default_stream(); + auto const int32_dtype = cudf::data_type{cudf::type_id::INT32}; + constexpr int32_t kExpected = 42; + + // Build a flat n_fields-field object with 1-byte value offsets where field `target_fid` + // holds INT32(kExpected) and all others hold BOOLEAN_TRUE. + auto build_flat = [&](int n_fields, int target_fid) { + auto const payload = enc_int32(kExpected); + constexpr uint8_t kBoolTrue = 0x04; + std::vector val{make_variant_object_header(), static_cast(n_fields)}; + for (int i = 0; i < n_fields; ++i) + val.push_back(static_cast(i)); + uint8_t off = 0; + for (int i = 0; i < n_fields; ++i) { + val.push_back(off); + off = static_cast(off + (i == target_fid ? payload.size() : 1u)); + } + val.push_back(off); + for (int i = 0; i < n_fields; ++i) { + if (i == target_fid) { + val.insert(val.end(), payload.begin(), payload.end()); + } else { + val.push_back(kBoolTrue); + } + } + return val; + }; + + // Case 1: total == 255 (85 keys × 3 bytes). Stays at 1-byte offsets. + // Extract the first key "k00" (field ID 0). + { + SCOPED_TRACE("total=255, 1-byte offsets"); + auto const keys = make_numeric_keys(85); + auto col = wrap_single_variant(build_metadata(keys), build_flat(85, /*target_fid=*/0)); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, "k00", int32_dtype, stream); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, + cudf::test::fixed_width_column_wrapper{kExpected}); + } + + // Case 2: total == 256 (84 keys × 3 bytes + "long" at 4 bytes). Switches to 2-byte offsets. + // "long" sorts after all "kXX" keys ('l' > 'k'), so it becomes field ID 84. + { + SCOPED_TRACE("total=256, 2-byte offsets"); + auto keys = make_numeric_keys(84); + keys.emplace_back("long"); + auto col = wrap_single_variant(build_metadata(keys), build_flat(85, /*target_fid=*/84)); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, "long", int32_dtype, stream); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, + cudf::test::fixed_width_column_wrapper{kExpected}); + } +} + TEST_F(ExtractVariantFieldTest, MalformedVariantDataYieldsNull) { // The column shape is a valid STRUCT, list>, but the VARIANT bytes are From f175243165bf0c2186613386519dc80b50dc70b7 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 01:28:25 +0000 Subject: [PATCH 06/12] review changes --- .../parquet/experimental/variant/extract.cpp | 238 +++++++++++------- .../io/experimental/variant_extract_test.cpp | 98 ++++---- 2 files changed, 205 insertions(+), 131 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index a13e83cd2969..38212d3ffc4b 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -7,6 +7,7 @@ #include #include #include +#include #include #include #include @@ -19,12 +20,61 @@ #include #include #include +#include #include #include #include namespace { +using cudf::io::parquet::experimental::variant_basic_type; +using cudf::io::parquet::experimental::variant_primitive_type; + +// The leaf value type exercised by a benchmark row, selected via the nvbench "type" string axis. +enum class bench_variant_type : uint8_t { INT32, FLOAT, BOOL, STRING, ARRAY }; + +// Parse the nvbench "type" string axis value into a bench_variant_type. +bench_variant_type parse_bench_variant_type(std::string const& type_str) +{ + if (type_str == "int32_t") return bench_variant_type::INT32; + if (type_str == "float") return bench_variant_type::FLOAT; + if (type_str == "bool") return bench_variant_type::BOOL; + if (type_str == "string") return bench_variant_type::STRING; + if (type_str == "array") return bench_variant_type::ARRAY; + CUDF_FAIL("Unrecognized benchmark type: " + type_str); +} + +// Compose a value-metadata header byte from a basic type and its 6-bit value_header. +// See cpp/tests/io/experimental/variant_extract_test.cpp for the header byte layout. +constexpr uint8_t make_variant_header(variant_basic_type basic, uint8_t value_header) +{ + return static_cast(static_cast(basic) | (value_header << 2)); +} + +// Header byte for a primitive value of the given physical type. +constexpr uint8_t make_variant_primitive(variant_primitive_type type) +{ + return make_variant_header(variant_basic_type::PRIMITIVE, static_cast(type)); +} + +// Header byte for a short string of the given length (must fit in 6 bits: 0..63). +constexpr uint8_t make_variant_short_string_header(std::size_t length) +{ + return make_variant_header(variant_basic_type::SHORT_STRING, static_cast(length)); +} + +// Header byte for an object value with 1-byte field ids and 1-byte offsets (value_header == 0). +constexpr uint8_t make_variant_object_header() +{ + return make_variant_header(variant_basic_type::OBJECT, 0); +} + +// Header byte for an array value with 1-byte count and 1-byte offsets (value_header == 0). +constexpr uint8_t make_variant_array_header() +{ + return make_variant_header(variant_basic_type::ARRAY, 0); +} + void append_le(std::vector& out, uint64_t bits, int width) { for (int i = 0; i < width; ++i) { @@ -64,65 +114,62 @@ std::vector build_metadata(std::vector const& keys) } // Wrap `inner` as the sole field (field id `fid`) of a 1-field VARIANT object. -// Uses 1-byte field_id_size and 1-byte field_offset_size (value_header=0 → header=0x02). -std::vector wrap_in_object(uint8_t fid, std::vector const& inner) +// Uses 1-byte field_id_size and 1-byte field_offset_size (value_header=0). +std::vector wrap_in_object(uint8_t fid, std::span inner) { // Format: object_header(1) + num_fields(1) + fid(1) + offset[0]=0(1) + offset[1]=size(1) + data - std::vector out{0x02, 0x01, fid, 0x00, static_cast(inner.size())}; + std::vector out{ + make_variant_object_header(), 0x01, fid, 0x00, static_cast(inner.size())}; out.insert(out.end(), inner.begin(), inner.end()); return out; } // Build the leaf VARIANT value blob for the requested type. -// -// Header byte composition: (physical_type_id << 2) | basic_type -// PRIMITIVE basic_type = 0, so header = physical_type_id << 2 -// SHORT_STRING basic_type = 1, so header = (length << 2) | 1 -// ARRAY basic_type = 3, so header = (value_header << 2) | 3 -// -// Physical type IDs used: -// INT32 = 5 → header 0x14 -// FLOAT32 = 14 → header 0x38 -// BOOL_TRUE= 1 → header 0x04 -std::vector build_leaf_value(std::string const& type_str) +std::vector build_leaf_value(bench_variant_type type) { - if (type_str == "int32_t") { - std::vector out{0x14}; - append_le(out, 42u, 4); - return out; - } - if (type_str == "float") { - std::vector out{0x38}; - float const f = 1.0f; - uint32_t u; - std::memcpy(&u, &f, 4); - append_le(out, u, 4); - return out; - } - if (type_str == "bool") { - return {0x04}; // BOOLEAN_TRUE - } - if (type_str == "string") { - // Short string "hello" (5 bytes): (5 << 2) | 1 = 0x15 - return {0x15, 'h', 'e', 'l', 'l', 'o'}; + switch (type) { + case bench_variant_type::INT32: { + std::vector out{make_variant_primitive(variant_primitive_type::INT32)}; + append_le(out, 42u, 4); + return out; + } + case bench_variant_type::FLOAT: { + std::vector out{make_variant_primitive(variant_primitive_type::FLOAT32)}; + float const f = 1.0f; + uint32_t u; + std::memcpy(&u, &f, 4); + append_le(out, u, 4); + return out; + } + case bench_variant_type::BOOL: + return {make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)}; + case bench_variant_type::STRING: { + // Short string "hello" (5 bytes). + auto const s = std::string{"hello"}; + std::vector out{make_variant_short_string_header(s.size())}; + out.insert(out.end(), s.begin(), s.end()); + return out; + } + case bench_variant_type::ARRAY: { + // VARIANT array of two INT32 values [42, 99]; element [1] is accessed in the benchmark. + // 2 elements, offsets [0, 5, 10], then INT32(42) and INT32(99) (5 bytes each). + std::vector out{make_variant_array_header(), 0x02, 0x00, 0x05, 0x0a}; + out.push_back(make_variant_primitive(variant_primitive_type::INT32)); + append_le(out, 42u, 4); + out.push_back(make_variant_primitive(variant_primitive_type::INT32)); + append_le(out, 99u, 4); + return out; + } + default: CUDF_FAIL("Unsupported benchmark leaf type"); } - // "array": VARIANT array of two INT32 values [42, 99]; element [1] is accessed in the benchmark. - // Array header 0x03: basic_type=ARRAY(3), value_header=0 (1-byte count, 1-byte offsets). - // 2 elements, offsets [0, 5, 10], then INT32(42) and INT32(99) (5 bytes each). - std::vector out{0x03, 0x02, 0x00, 0x05, 0x0a}; - out.push_back(0x14); - append_le(out, 42u, 4); - out.push_back(0x14); - append_le(out, 99u, 4); - return out; } // Build the full hit-row value blob by wrapping the leaf in `nesting` object levels. // Keys a,b,c,d,e map to field IDs 0,1,2,3,4 in the shared dictionary. // For path a.b.c.d.e the outermost object uses fid=0 ("a"). -std::vector build_hit_value(std::string const& type_str, int nesting) +std::vector build_hit_value(bench_variant_type type, int nesting) { - auto val = build_leaf_value(type_str); + auto val = build_leaf_value(type); for (int i = nesting - 1; i >= 0; --i) { val = wrap_in_object(static_cast(i), val); } @@ -134,32 +181,44 @@ std::vector build_hit_value(std::string const& type_str, int nesting) // dictionary), so traversal fails at the first key lookup while the row remains non-null. // For cast_variant rows (nesting=0, non-array): a different primitive type so the cast returns // null. -std::vector build_miss_value(int nesting, bool is_array, std::string const& type_str) +std::vector build_miss_value(int nesting, bool is_array, bench_variant_type type) { if (nesting == 0 && !is_array) { // Wrong-type primitive for the cast path. - if (type_str == "bool") { - std::vector out{0x14}; - append_le(out, 0u, 4); - return out; + switch (type) { + case bench_variant_type::BOOL: { + std::vector out{make_variant_primitive(variant_primitive_type::INT32)}; + append_le(out, 0u, 4); + return out; + } + default: return {make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)}; } - return {0x04}; // BOOLEAN_TRUE } // "z" is always the last key in the dictionary, at field ID = nesting. - return wrap_in_object(static_cast(nesting), build_leaf_value(type_str)); + return wrap_in_object(static_cast(nesting), build_leaf_value(type)); } -// Build a VARIANT struct column (STRUCT, list>) from per-row byte vectors. -std::unique_ptr build_variant_column( - std::vector> const& meta_rows, - std::vector> const& val_rows, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) +// Build a std::span view over each row of a vector-of-byte-vectors, without copying row data. +std::vector> to_spans(std::vector> const& rows) +{ + std::vector> spans; + spans.reserve(rows.size()); + for (auto const& row : rows) { + spans.emplace_back(row); + } + return spans; +} + +// Build a VARIANT struct column (STRUCT, list>) from per-row byte spans. +std::unique_ptr build_variant_column(std::span> meta_rows, + std::span> val_rows, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) { auto const n = static_cast(meta_rows.size()); auto build_list_col = - [&](std::vector> const& rows) -> std::unique_ptr { + [&](std::span> rows) -> std::unique_ptr { std::vector offsets(n + 1, 0); std::vector flat; for (cudf::size_type i = 0; i < n; ++i) { @@ -215,13 +274,13 @@ std::vector get_dict_keys_for_fields(int num_fields) } // Build a flat object with `num_fields` fields using 1-byte field IDs and 1-byte offsets. -// Field `target_fid` holds `inner`; all other fields hold a dummy BOOLEAN_TRUE (0x04). +// Field `target_fid` holds `inner`; all other fields hold a dummy BOOLEAN_TRUE. std::vector build_flat_object(int num_fields, int target_fid, - std::vector const& inner) + std::span inner) { // object_header(1) + num_fields(1) + field_ids(num_fields) + offsets(num_fields+1) + data - std::vector out{0x02, static_cast(num_fields)}; + std::vector out{make_variant_object_header(), static_cast(num_fields)}; for (int i = 0; i < num_fields; ++i) { out.push_back(static_cast(i)); } @@ -235,7 +294,7 @@ std::vector build_flat_object(int num_fields, if (i == target_fid) { out.insert(out.end(), inner.begin(), inner.end()); } else { - out.push_back(0x04); // BOOLEAN_TRUE dummy + out.push_back(make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)); // dummy } } return out; @@ -256,13 +315,17 @@ std::string get_path(int nesting, bool is_array) return path; } -cudf::data_type get_target_type(std::string const& type_str) +cudf::data_type get_target_type(bench_variant_type type) { - if (type_str == "float") return cudf::data_type{cudf::type_id::FLOAT32}; - if (type_str == "bool") return cudf::data_type{cudf::type_id::BOOL8}; - if (type_str == "string") return cudf::data_type{cudf::type_id::STRING}; - // "int32_t" and "array" (element access yields INT32) - return cudf::data_type{cudf::type_id::INT32}; + switch (type) { + case bench_variant_type::FLOAT: return cudf::data_type{cudf::type_id::FLOAT32}; + case bench_variant_type::BOOL: return cudf::data_type{cudf::type_id::BOOL8}; + case bench_variant_type::STRING: return cudf::data_type{cudf::type_id::STRING}; + // "array": element access yields INT32. + case bench_variant_type::INT32: + case bench_variant_type::ARRAY: return cudf::data_type{cudf::type_id::INT32}; + default: CUDF_FAIL("Unsupported benchmark target type"); + } } } // namespace @@ -270,14 +333,15 @@ cudf::data_type get_target_type(std::string const& type_str) // Assign each row randomly as a hit or miss rather than using contiguous strided ranges, // so the memory access pattern doesn't accidentally favour cache locality. void fill_val_rows(std::vector>& val_rows, - std::vector const& hit_val, - std::vector const& miss_val, + std::span hit_val, + std::span miss_val, int hit_rate) { std::mt19937 rng{42}; std::uniform_int_distribution dist{0, 99}; for (auto& row : val_rows) { - row = (dist(rng) < hit_rate) ? hit_val : miss_val; + auto const& src = (dist(rng) < hit_rate) ? hit_val : miss_val; + row.assign(src.begin(), src.end()); } } @@ -288,21 +352,23 @@ static void bench_variant_cast(nvbench::state& state) auto mr = cudf::get_current_device_resource_ref(); auto const num_rows = static_cast(state.get_int64("num_rows")); - auto const type_str = state.get_string("type"); + auto const type = parse_bench_variant_type(state.get_string("type")); auto const hit_rate = static_cast(state.get_int64("hit_rate")); auto const meta_blob = build_metadata(get_dict_keys(0)); - auto const hit_val = build_leaf_value(type_str); - auto const miss_val = build_miss_value(0, /*is_array=*/false, type_str); + auto const hit_val = build_leaf_value(type); + auto const miss_val = build_miss_value(0, /*is_array=*/false, type); std::vector> meta_rows(num_rows, meta_blob); std::vector> val_rows(num_rows); fill_val_rows(val_rows, hit_val, miss_val, hit_rate); - auto col = build_variant_column(meta_rows, val_rows, stream, mr); + auto meta_spans = to_spans(meta_rows); + auto val_spans = to_spans(val_rows); + auto col = build_variant_column(meta_spans, val_spans, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); - auto const target_type = get_target_type(type_str); + auto const target_type = get_target_type(type); state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { @@ -324,24 +390,26 @@ static void bench_variant_extract_nesting(nvbench::state& state) auto mr = cudf::get_current_device_resource_ref(); auto const num_rows = static_cast(state.get_int64("num_rows")); - auto const type_str = state.get_string("type"); + auto const type = parse_bench_variant_type(state.get_string("type")); auto const nesting = static_cast(state.get_int64("nesting")); auto const hit_rate = static_cast(state.get_int64("hit_rate")); - bool const is_array = (type_str == "array"); + bool const is_array = (type == bench_variant_type::ARRAY); auto const meta_blob = build_metadata(get_dict_keys(nesting)); - auto const hit_val = build_hit_value(type_str, nesting); - auto const miss_val = build_miss_value(nesting, is_array, type_str); + auto const hit_val = build_hit_value(type, nesting); + auto const miss_val = build_miss_value(nesting, is_array, type); std::vector> meta_rows(num_rows, meta_blob); std::vector> val_rows(num_rows); fill_val_rows(val_rows, hit_val, miss_val, hit_rate); - auto col = build_variant_column(meta_rows, val_rows, stream, mr); + auto meta_spans = to_spans(meta_rows); + auto val_spans = to_spans(val_rows); + auto col = build_variant_column(meta_spans, val_spans, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); - auto const target_type = get_target_type(type_str); + auto const target_type = get_target_type(type); auto const path = get_path(nesting, is_array); state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); @@ -374,7 +442,7 @@ static void bench_variant_extract_fields(nvbench::state& state) int const target_fid = (field_pos_str == "last") ? (num_fields - 1) : 0; auto const meta_blob = build_metadata(get_dict_keys_for_fields(num_fields)); - auto const leaf = build_leaf_value("int32_t"); + auto const leaf = build_leaf_value(bench_variant_type::INT32); auto const hit_val = build_flat_object(num_fields, target_fid, leaf); // Miss: object keyed on "z" (field ID = num_fields), so the lookup fails. auto const miss_val = wrap_in_object(static_cast(num_fields), leaf); @@ -383,7 +451,9 @@ static void bench_variant_extract_fields(nvbench::state& state) std::vector> val_rows(num_rows); fill_val_rows(val_rows, hit_val, miss_val, hit_rate); - auto col = build_variant_column(meta_rows, val_rows, stream, mr); + auto meta_spans = to_spans(meta_rows); + auto val_spans = to_spans(val_rows); + auto col = build_variant_column(meta_spans, val_spans, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); std::string const path = diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 69ae171ccaaa..fcebe94c98d6 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -540,28 +540,33 @@ inline cudf::test::structs_column_wrapper wrap_multi_row_variant( // Header bits [7:6] = offset_size_minus_one; bits [3:0] = version (1). inline std::vector build_metadata(std::vector const& keys) { - uint32_t total = 0; - for (auto const& k : keys) - total += static_cast(k.size()); + constexpr uint8_t kVariantMetadataVersion = 0x01; + constexpr int kMetadataOffsetSizeShift = 6; + constexpr uint32_t kMaxSingleByteOffsetSum = 255u; - int const offset_size = (total > 255u) ? 2 : 1; - std::vector out{static_cast(0x01 | ((offset_size - 1) << 6))}; + uint32_t total_key_bytes = 0; + for (auto const& key : keys) + total_key_bytes += static_cast(key.size()); - auto write_le = [&](uint32_t v) { - for (int i = 0; i < offset_size; ++i) - out.push_back(static_cast(v >> (8 * i))); + int const offset_size = (total_key_bytes > kMaxSingleByteOffsetSum) ? 2 : 1; + std::vector out{static_cast(kVariantMetadataVersion | + ((offset_size - 1) << kMetadataOffsetSizeShift))}; + + auto write_little_endian_offset = [&](uint32_t value) { + for (int byte_index = 0; byte_index < offset_size; ++byte_index) + out.push_back(static_cast(value >> (8 * byte_index))); }; - write_le(static_cast(keys.size())); + write_little_endian_offset(static_cast(keys.size())); - uint32_t running = 0; - write_le(0u); - for (auto const& k : keys) { - running += static_cast(k.size()); - write_le(running); + uint32_t running_offset = 0; + write_little_endian_offset(0u); + for (auto const& key : keys) { + running_offset += static_cast(key.size()); + write_little_endian_offset(running_offset); } - for (auto const& k : keys) { - out.insert(out.end(), k.begin(), k.end()); + for (auto const& key : keys) { + out.insert(out.end(), key.begin(), key.end()); } return out; } @@ -789,25 +794,25 @@ TEST_F(ExtractVariantFieldTest, LargeDictionary100FieldsExtractLast) auto const keys = make_numeric_keys(100); auto const meta = build_metadata(keys); - constexpr int n_fields = 100; - constexpr int target_fid = 99; - auto const target_val = enc_int32(target_fid); - constexpr uint8_t bool_true_byte = 0x04; - - std::vector val{make_variant_object_header(), static_cast(n_fields)}; - for (int i = 0; i < n_fields; ++i) - val.push_back(static_cast(i)); // field IDs - uint8_t off = 0; - for (int i = 0; i < n_fields; ++i) { - val.push_back(off); - off = static_cast(off + (i == target_fid ? target_val.size() : 1u)); + constexpr int field_count = 100; + constexpr int target_fid = 99; + auto const target_val = enc_int32(target_fid); + + std::vector val{make_variant_object_header(), static_cast(field_count)}; + for (int fid = 0; fid < field_count; ++fid) + val.push_back(static_cast(fid)); + uint8_t field_offset = 0; + for (int fid = 0; fid < field_count; ++fid) { + val.push_back(field_offset); + field_offset = + static_cast(field_offset + (fid == target_fid ? target_val.size() : 1u)); } - val.push_back(off); // sentinel offset after last field - for (int i = 0; i < n_fields; ++i) { - if (i == target_fid) { + val.push_back(field_offset); + for (int fid = 0; fid < field_count; ++fid) { + if (fid == target_fid) { val.insert(val.end(), target_val.begin(), target_val.end()); } else { - val.push_back(bool_true_byte); + val.push_back(make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)); } } @@ -827,25 +832,24 @@ TEST_F(ExtractVariantFieldTest, MetadataOffsetSizeThresholdBoundary) auto const int32_dtype = cudf::data_type{cudf::type_id::INT32}; constexpr int32_t kExpected = 42; - // Build a flat n_fields-field object with 1-byte value offsets where field `target_fid` + // Build a flat field_count-field object with 1-byte value offsets where field `target_fid` // holds INT32(kExpected) and all others hold BOOLEAN_TRUE. - auto build_flat = [&](int n_fields, int target_fid) { - auto const payload = enc_int32(kExpected); - constexpr uint8_t kBoolTrue = 0x04; - std::vector val{make_variant_object_header(), static_cast(n_fields)}; - for (int i = 0; i < n_fields; ++i) - val.push_back(static_cast(i)); - uint8_t off = 0; - for (int i = 0; i < n_fields; ++i) { - val.push_back(off); - off = static_cast(off + (i == target_fid ? payload.size() : 1u)); + auto build_flat = [&](int field_count, int target_fid) { + auto const payload = enc_int32(kExpected); + std::vector val{make_variant_object_header(), static_cast(field_count)}; + for (int fid = 0; fid < field_count; ++fid) + val.push_back(static_cast(fid)); + uint8_t field_offset = 0; + for (int fid = 0; fid < field_count; ++fid) { + val.push_back(field_offset); + field_offset = static_cast(field_offset + (fid == target_fid ? payload.size() : 1u)); } - val.push_back(off); - for (int i = 0; i < n_fields; ++i) { - if (i == target_fid) { + val.push_back(field_offset); + for (int fid = 0; fid < field_count; ++fid) { + if (fid == target_fid) { val.insert(val.end(), payload.begin(), payload.end()); } else { - val.push_back(kBoolTrue); + val.push_back(make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)); } } return val; From e6bcd531b598a041b5dbf09e9dad315224dc3f3e Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 01:33:17 +0000 Subject: [PATCH 07/12] finishing bdice comments --- .../parquet/experimental/variant/extract.cpp | 20 ++++++++++--------- .../io/experimental/variant_extract_test.cpp | 12 +++++++---- 2 files changed, 19 insertions(+), 13 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 38212d3ffc4b..c2bf70a3c95c 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -36,11 +36,11 @@ enum class bench_variant_type : uint8_t { INT32, FLOAT, BOOL, STRING, ARRAY }; // Parse the nvbench "type" string axis value into a bench_variant_type. bench_variant_type parse_bench_variant_type(std::string const& type_str) { - if (type_str == "int32_t") return bench_variant_type::INT32; - if (type_str == "float") return bench_variant_type::FLOAT; - if (type_str == "bool") return bench_variant_type::BOOL; - if (type_str == "string") return bench_variant_type::STRING; - if (type_str == "array") return bench_variant_type::ARRAY; + if (type_str == "int32_t") { return bench_variant_type::INT32; } + if (type_str == "float") { return bench_variant_type::FLOAT; } + if (type_str == "bool") { return bench_variant_type::BOOL; } + if (type_str == "string") { return bench_variant_type::STRING; } + if (type_str == "array") { return bench_variant_type::ARRAY; } CUDF_FAIL("Unrecognized benchmark type: " + type_str); } @@ -88,15 +88,17 @@ void append_le(std::vector& out, uint64_t bits, int width) std::vector build_metadata(std::vector const& keys) { uint32_t total = 0; - for (auto const& k : keys) + for (auto const& k : keys) { total += static_cast(k.size()); + } int const offset_size = (total > 255u) ? 2 : 1; std::vector out{static_cast(0x01 | ((offset_size - 1) << 6))}; auto write_le = [&](uint32_t v) { - for (int i = 0; i < offset_size; ++i) + for (int i = 0; i < offset_size; ++i) { out.push_back(static_cast(v >> (8 * i))); + } }; write_le(static_cast(keys.size())); @@ -308,10 +310,10 @@ std::string get_path(int nesting, bool is_array) { std::string path; for (int i = 0; i < nesting; ++i) { - if (i > 0) path += '.'; + if (i > 0) { path += '.'; } path += static_cast('a' + i); } - if (is_array) path += "[1]"; + if (is_array) { path += "[1]"; } return path; } diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index fcebe94c98d6..70face78ad61 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -545,16 +545,18 @@ inline std::vector build_metadata(std::vector const& keys) constexpr uint32_t kMaxSingleByteOffsetSum = 255u; uint32_t total_key_bytes = 0; - for (auto const& key : keys) + for (auto const& key : keys) { total_key_bytes += static_cast(key.size()); + } int const offset_size = (total_key_bytes > kMaxSingleByteOffsetSum) ? 2 : 1; std::vector out{static_cast(kVariantMetadataVersion | ((offset_size - 1) << kMetadataOffsetSizeShift))}; auto write_little_endian_offset = [&](uint32_t value) { - for (int byte_index = 0; byte_index < offset_size; ++byte_index) + for (int byte_index = 0; byte_index < offset_size; ++byte_index) { out.push_back(static_cast(value >> (8 * byte_index))); + } }; write_little_endian_offset(static_cast(keys.size())); @@ -799,8 +801,9 @@ TEST_F(ExtractVariantFieldTest, LargeDictionary100FieldsExtractLast) auto const target_val = enc_int32(target_fid); std::vector val{make_variant_object_header(), static_cast(field_count)}; - for (int fid = 0; fid < field_count; ++fid) + for (int fid = 0; fid < field_count; ++fid) { val.push_back(static_cast(fid)); + } uint8_t field_offset = 0; for (int fid = 0; fid < field_count; ++fid) { val.push_back(field_offset); @@ -837,8 +840,9 @@ TEST_F(ExtractVariantFieldTest, MetadataOffsetSizeThresholdBoundary) auto build_flat = [&](int field_count, int target_fid) { auto const payload = enc_int32(kExpected); std::vector val{make_variant_object_header(), static_cast(field_count)}; - for (int fid = 0; fid < field_count; ++fid) + for (int fid = 0; fid < field_count; ++fid) { val.push_back(static_cast(fid)); + } uint8_t field_offset = 0; for (int fid = 0; fid < field_count; ++fid) { val.push_back(field_offset); From a4a49b7373f60528b828ccdb17d00d9b17310364 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 02:11:30 +0000 Subject: [PATCH 08/12] vukasin's changes --- .../parquet/experimental/variant/extract.cpp | 185 +++++++++++------- 1 file changed, 109 insertions(+), 76 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index c2bf70a3c95c..68bfee9ca226 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -3,6 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include + #include #include #include @@ -17,8 +19,10 @@ #include +#include #include #include +#include #include #include #include @@ -82,35 +86,44 @@ void append_le(std::vector& out, uint64_t bits, int width) } } -// Build a V1 VARIANT metadata blob for a sorted key dictionary. +// Build a V1 VARIANT metadata blob for a sorted key dictionary. Callers must pass `keys` in +// ascending sorted order; the sorted-strings header bit is always set to reflect that. // Uses 2-byte offsets when the total string length exceeds 255 bytes; 1-byte otherwise. -// Header bits [7:6] = offset_size_minus_one; bits [3:0] = version (1). +// Header bits [7:6] = offset_size_minus_one; bit [4] = sorted_strings; bits [3:0] = version (1). std::vector build_metadata(std::vector const& keys) { - uint32_t total = 0; - for (auto const& k : keys) { - total += static_cast(k.size()); + constexpr uint8_t kVariantMetadataVersion = 0x01; + constexpr uint8_t kVariantMetadataSorted = 0x10; + constexpr int kMetadataOffsetSizeShift = 6; + constexpr uint32_t kMaxSingleByteOffsetSum = 255u; + + uint32_t total_key_bytes = 0; + for (auto const& key : keys) { + total_key_bytes += static_cast(key.size()); } - int const offset_size = (total > 255u) ? 2 : 1; - std::vector out{static_cast(0x01 | ((offset_size - 1) << 6))}; + int const offset_size = (total_key_bytes > kMaxSingleByteOffsetSum) ? 2 : 1; + std::vector out{static_cast(kVariantMetadataVersion | kVariantMetadataSorted | + ((offset_size - 1) << kMetadataOffsetSizeShift))}; + out.reserve(out.size() + static_cast(offset_size) * (keys.size() + 2) + + total_key_bytes); - auto write_le = [&](uint32_t v) { - for (int i = 0; i < offset_size; ++i) { - out.push_back(static_cast(v >> (8 * i))); + auto write_little_endian_offset = [&](uint32_t value) { + for (int byte_index = 0; byte_index < offset_size; ++byte_index) { + out.push_back(static_cast(value >> (8 * byte_index))); } }; - write_le(static_cast(keys.size())); + write_little_endian_offset(static_cast(keys.size())); - uint32_t running = 0; - write_le(0u); - for (auto const& k : keys) { - running += static_cast(k.size()); - write_le(running); + uint32_t running_offset = 0; + write_little_endian_offset(0u); + for (auto const& key : keys) { + running_offset += static_cast(key.size()); + write_little_endian_offset(running_offset); } - for (auto const& k : keys) { - out.insert(out.end(), k.begin(), k.end()); + for (auto const& key : keys) { + out.insert(out.end(), key.begin(), key.end()); } return out; } @@ -179,7 +192,7 @@ std::vector build_hit_value(bench_variant_type type, int nesting) } // Build the miss-row value blob: a valid VARIANT that won't match the target path or type. -// For extract_variant_field rows: a 1-level object keyed on "z" (field ID = nesting in the +// For get_variant_field rows: a 1-level object keyed on "z" (field ID = nesting in the // dictionary), so traversal fails at the first key lookup while the row remains non-null. // For cast_variant rows (nesting=0, non-array): a different primitive type so the cast returns // null. @@ -200,15 +213,14 @@ std::vector build_miss_value(int nesting, bool is_array, bench_variant_ return wrap_in_object(static_cast(nesting), build_leaf_value(type)); } -// Build a std::span view over each row of a vector-of-byte-vectors, without copying row data. -std::vector> to_spans(std::vector> const& rows) +// Zero-pad the shorter of `hit_val`/`miss_val` so both end up the same length. VARIANT decoders +// only ever read the bytes their own header/offsets describe, so trailing padding is inert; this +// keeps a row's size from being a confound for hit vs. miss access-pattern benchmarking. +void pad_to_equal_size(std::vector& hit_val, std::vector& miss_val) { - std::vector> spans; - spans.reserve(rows.size()); - for (auto const& row : rows) { - spans.emplace_back(row); - } - return spans; + auto const target_size = std::max(hit_val.size(), miss_val.size()); + hit_val.resize(target_size, uint8_t{0}); + miss_val.resize(target_size, uint8_t{0}); } // Build a VARIANT struct column (STRUCT, list>) from per-row byte spans. @@ -222,7 +234,12 @@ std::unique_ptr build_variant_column(std::span> rows) -> std::unique_ptr { std::vector offsets(n + 1, 0); + auto const total_bytes = std::accumulate( + rows.begin(), rows.end(), std::size_t{0}, [](std::size_t acc, auto const& row) { + return acc + row.size(); + }); std::vector flat; + flat.reserve(total_bytes); for (cudf::size_type i = 0; i < n; ++i) { flat.insert(flat.end(), rows[i].begin(), rows[i].end()); offsets[i + 1] = static_cast(flat.size()); @@ -283,6 +300,7 @@ std::vector build_flat_object(int num_fields, { // object_header(1) + num_fields(1) + field_ids(num_fields) + offsets(num_fields+1) + data std::vector out{make_variant_object_header(), static_cast(num_fields)}; + out.reserve(out.size() + static_cast(3 * num_fields) + inner.size()); for (int i = 0; i < num_fields; ++i) { out.push_back(static_cast(i)); } @@ -330,23 +348,26 @@ cudf::data_type get_target_type(bench_variant_type type) } } -} // namespace - -// Assign each row randomly as a hit or miss rather than using contiguous strided ranges, -// so the memory access pattern doesn't accidentally favour cache locality. -void fill_val_rows(std::vector>& val_rows, - std::span hit_val, - std::span miss_val, - int hit_rate) +// Assign each row randomly as a hit or miss rather than using contiguous strided ranges, so the +// memory access pattern doesn't accidentally favour cache locality. Rows are spans aliasing +// `hit_val`/`miss_val` directly, avoiding a per-row byte copy. +std::vector> fill_val_rows(cudf::size_type num_rows, + std::span hit_val, + std::span miss_val, + int hit_rate) { std::mt19937 rng{42}; std::uniform_int_distribution dist{0, 99}; - for (auto& row : val_rows) { - auto const& src = (dist(rng) < hit_rate) ? hit_val : miss_val; - row.assign(src.begin(), src.end()); + std::vector> val_rows; + val_rows.reserve(num_rows); + for (cudf::size_type i = 0; i < num_rows; ++i) { + val_rows.push_back((dist(rng) < hit_rate) ? hit_val : miss_val); } + return val_rows; } +} // namespace + // Benchmarks cast_variant: each row's value IS the leaf primitive (no path traversal). static void bench_variant_cast(nvbench::state& state) { @@ -358,25 +379,30 @@ static void bench_variant_cast(nvbench::state& state) auto const hit_rate = static_cast(state.get_int64("hit_rate")); auto const meta_blob = build_metadata(get_dict_keys(0)); - auto const hit_val = build_leaf_value(type); - auto const miss_val = build_miss_value(0, /*is_array=*/false, type); - - std::vector> meta_rows(num_rows, meta_blob); - std::vector> val_rows(num_rows); - fill_val_rows(val_rows, hit_val, miss_val, hit_rate); + auto hit_val = build_leaf_value(type); + auto miss_val = build_miss_value(0, /*is_array=*/false, type); + pad_to_equal_size(hit_val, miss_val); - auto meta_spans = to_spans(meta_rows); - auto val_spans = to_spans(val_rows); - auto col = build_variant_column(meta_spans, val_spans, stream, mr); + std::vector> meta_spans(num_rows, std::span{meta_blob}); + auto val_spans = fill_val_rows(num_rows, hit_val, miss_val, hit_rate); + auto col = build_variant_column(meta_spans, val_spans, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); auto const target_type = get_target_type(type); + auto const data_size = static_cast(num_rows) * (meta_blob.size() + hit_val.size()); + auto mem_stats_logger = cudf::memory_stats_logger(); + mr = cudf::get_current_device_resource_ref(); state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { std::ignore = cudf::io::parquet::experimental::cast_variant(col->view().child(1), target_type, stream, mr); }); + + auto const time = state.get_summary("nv/cold/time/gpu/mean").get_float64("value"); + state.add_element_count(static_cast(data_size) / time, "bytes_per_second"); + state.add_buffer_size( + mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); } NVBENCH_BENCH(bench_variant_cast) @@ -385,7 +411,8 @@ NVBENCH_BENCH(bench_variant_cast) .add_string_axis("type", {"string", "float", "bool", "int32_t"}) .add_int64_axis("hit_rate", {20, 80}); -// Benchmarks extract_variant_field with varying path depth (nesting >= 1). +// Benchmarks get_variant_field with varying path depth (nesting >= 1). Casting is exercised +// separately by bench_variant_cast, so this isolates pure path-traversal cost. static void bench_variant_extract_nesting(nvbench::state& state) { auto stream = cudf::get_default_stream(); @@ -399,26 +426,29 @@ static void bench_variant_extract_nesting(nvbench::state& state) bool const is_array = (type == bench_variant_type::ARRAY); auto const meta_blob = build_metadata(get_dict_keys(nesting)); - auto const hit_val = build_hit_value(type, nesting); - auto const miss_val = build_miss_value(nesting, is_array, type); - - std::vector> meta_rows(num_rows, meta_blob); - std::vector> val_rows(num_rows); - fill_val_rows(val_rows, hit_val, miss_val, hit_rate); + auto hit_val = build_hit_value(type, nesting); + auto miss_val = build_miss_value(nesting, is_array, type); + pad_to_equal_size(hit_val, miss_val); - auto meta_spans = to_spans(meta_rows); - auto val_spans = to_spans(val_rows); - auto col = build_variant_column(meta_spans, val_spans, stream, mr); + std::vector> meta_spans(num_rows, std::span{meta_blob}); + auto val_spans = fill_val_rows(num_rows, hit_val, miss_val, hit_rate); + auto col = build_variant_column(meta_spans, val_spans, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); - auto const target_type = get_target_type(type); - auto const path = get_path(nesting, is_array); + auto const path = get_path(nesting, is_array); + auto const data_size = static_cast(num_rows) * (meta_blob.size() + hit_val.size()); + auto mem_stats_logger = cudf::memory_stats_logger(); + mr = cudf::get_current_device_resource_ref(); state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { - std::ignore = cudf::io::parquet::experimental::extract_variant_field( - col->view(), path, target_type, stream, mr); + std::ignore = cudf::io::parquet::experimental::get_variant_field(col->view(), path, stream, mr); }); + + auto const time = state.get_summary("nv/cold/time/gpu/mean").get_float64("value"); + state.add_element_count(static_cast(data_size) / time, "bytes_per_second"); + state.add_buffer_size( + mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); } NVBENCH_BENCH(bench_variant_extract_nesting) @@ -428,9 +458,9 @@ NVBENCH_BENCH(bench_variant_extract_nesting) .add_int64_axis("nesting", {1, 5}) .add_int64_axis("hit_rate", {20, 80}); -// Benchmarks extract_variant_field on a flat object, varying the total number of fields -// and whether the target field is first or last (probes binary search cost). -// Type is fixed to int32_t to isolate field-lookup overhead. +// Benchmarks get_variant_field on a flat object, varying the total number of fields and whether +// the target field is first or last (probes binary search cost). Type is fixed to int32_t to +// isolate field-lookup overhead; casting is exercised separately by bench_variant_cast. static void bench_variant_extract_fields(nvbench::state& state) { auto stream = cudf::get_default_stream(); @@ -445,28 +475,31 @@ static void bench_variant_extract_fields(nvbench::state& state) auto const meta_blob = build_metadata(get_dict_keys_for_fields(num_fields)); auto const leaf = build_leaf_value(bench_variant_type::INT32); - auto const hit_val = build_flat_object(num_fields, target_fid, leaf); + auto hit_val = build_flat_object(num_fields, target_fid, leaf); // Miss: object keyed on "z" (field ID = num_fields), so the lookup fails. - auto const miss_val = wrap_in_object(static_cast(num_fields), leaf); - - std::vector> meta_rows(num_rows, meta_blob); - std::vector> val_rows(num_rows); - fill_val_rows(val_rows, hit_val, miss_val, hit_rate); + auto miss_val = wrap_in_object(static_cast(num_fields), leaf); + pad_to_equal_size(hit_val, miss_val); - auto meta_spans = to_spans(meta_rows); - auto val_spans = to_spans(val_rows); - auto col = build_variant_column(meta_spans, val_spans, stream, mr); + std::vector> meta_spans(num_rows, std::span{meta_blob}); + auto val_spans = fill_val_rows(num_rows, hit_val, miss_val, hit_rate); + auto col = build_variant_column(meta_spans, val_spans, stream, mr); CUDF_CUDA_TRY(cudaStreamSynchronize(stream.value())); std::string const path = "f" + std::string(target_fid < 10 ? "0" : "") + std::to_string(target_fid); - auto const target_type = cudf::data_type{cudf::type_id::INT32}; + auto const data_size = static_cast(num_rows) * (meta_blob.size() + hit_val.size()); + auto mem_stats_logger = cudf::memory_stats_logger(); + mr = cudf::get_current_device_resource_ref(); state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch&) { - std::ignore = cudf::io::parquet::experimental::extract_variant_field( - col->view(), path, target_type, stream, mr); + std::ignore = cudf::io::parquet::experimental::get_variant_field(col->view(), path, stream, mr); }); + + auto const time = state.get_summary("nv/cold/time/gpu/mean").get_float64("value"); + state.add_element_count(static_cast(data_size) / time, "bytes_per_second"); + state.add_buffer_size( + mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage"); } NVBENCH_BENCH(bench_variant_extract_fields) From 6154685b17b46a0ccf0f570d8948166222b33aaf Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 21:26:45 +0000 Subject: [PATCH 09/12] subset of changes --- .../parquet/experimental/variant/extract.cpp | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 68bfee9ca226..784346f88b3d 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -56,7 +56,7 @@ constexpr uint8_t make_variant_header(variant_basic_type basic, uint8_t value_he } // Header byte for a primitive value of the given physical type. -constexpr uint8_t make_variant_primitive(variant_primitive_type type) +constexpr uint8_t make_variant_primitive_header(variant_primitive_type type) { return make_variant_header(variant_basic_type::PRIMITIVE, static_cast(type)); } @@ -144,12 +144,12 @@ std::vector build_leaf_value(bench_variant_type type) { switch (type) { case bench_variant_type::INT32: { - std::vector out{make_variant_primitive(variant_primitive_type::INT32)}; + std::vector out{make_variant_primitive_header(variant_primitive_type::INT32)}; append_le(out, 42u, 4); return out; } case bench_variant_type::FLOAT: { - std::vector out{make_variant_primitive(variant_primitive_type::FLOAT32)}; + std::vector out{make_variant_primitive_header(variant_primitive_type::FLOAT32)}; float const f = 1.0f; uint32_t u; std::memcpy(&u, &f, 4); @@ -157,7 +157,7 @@ std::vector build_leaf_value(bench_variant_type type) return out; } case bench_variant_type::BOOL: - return {make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)}; + return {make_variant_primitive_header(variant_primitive_type::BOOLEAN_TRUE)}; case bench_variant_type::STRING: { // Short string "hello" (5 bytes). auto const s = std::string{"hello"}; @@ -169,9 +169,9 @@ std::vector build_leaf_value(bench_variant_type type) // VARIANT array of two INT32 values [42, 99]; element [1] is accessed in the benchmark. // 2 elements, offsets [0, 5, 10], then INT32(42) and INT32(99) (5 bytes each). std::vector out{make_variant_array_header(), 0x02, 0x00, 0x05, 0x0a}; - out.push_back(make_variant_primitive(variant_primitive_type::INT32)); + out.push_back(make_variant_primitive_header(variant_primitive_type::INT32)); append_le(out, 42u, 4); - out.push_back(make_variant_primitive(variant_primitive_type::INT32)); + out.push_back(make_variant_primitive_header(variant_primitive_type::INT32)); append_le(out, 99u, 4); return out; } @@ -202,11 +202,11 @@ std::vector build_miss_value(int nesting, bool is_array, bench_variant_ // Wrong-type primitive for the cast path. switch (type) { case bench_variant_type::BOOL: { - std::vector out{make_variant_primitive(variant_primitive_type::INT32)}; + std::vector out{make_variant_primitive_header(variant_primitive_type::INT32)}; append_le(out, 0u, 4); return out; } - default: return {make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)}; + default: return {make_variant_primitive_header(variant_primitive_type::BOOLEAN_TRUE)}; } } // "z" is always the last key in the dictionary, at field ID = nesting. @@ -314,7 +314,7 @@ std::vector build_flat_object(int num_fields, if (i == target_fid) { out.insert(out.end(), inner.begin(), inner.end()); } else { - out.push_back(make_variant_primitive(variant_primitive_type::BOOLEAN_TRUE)); // dummy + out.push_back(make_variant_primitive_header(variant_primitive_type::BOOLEAN_TRUE)); // dummy } } return out; From 9ac95717fd31e7cecdac3e0a1e8ac8444709d581 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 16:26:54 -0500 Subject: [PATCH 10/12] Update cpp/benchmarks/io/parquet/experimental/variant/extract.cpp Co-authored-by: Vukasin Milovanovic --- cpp/benchmarks/io/parquet/experimental/variant/extract.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 68bfee9ca226..28291e4238e1 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -37,7 +37,6 @@ using cudf::io::parquet::experimental::variant_primitive_type; // The leaf value type exercised by a benchmark row, selected via the nvbench "type" string axis. enum class bench_variant_type : uint8_t { INT32, FLOAT, BOOL, STRING, ARRAY }; -// Parse the nvbench "type" string axis value into a bench_variant_type. bench_variant_type parse_bench_variant_type(std::string const& type_str) { if (type_str == "int32_t") { return bench_variant_type::INT32; } From 8f865e6d43edf3b374b9eca0e3772f7aed6bf04d Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 16:27:03 -0500 Subject: [PATCH 11/12] Update cpp/benchmarks/io/parquet/experimental/variant/extract.cpp Co-authored-by: Vukasin Milovanovic --- cpp/benchmarks/io/parquet/experimental/variant/extract.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 28291e4238e1..913c90d8ae8b 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -54,7 +54,6 @@ constexpr uint8_t make_variant_header(variant_basic_type basic, uint8_t value_he return static_cast(static_cast(basic) | (value_header << 2)); } -// Header byte for a primitive value of the given physical type. constexpr uint8_t make_variant_primitive(variant_primitive_type type) { return make_variant_header(variant_basic_type::PRIMITIVE, static_cast(type)); From 209fbcd0ce7eae3a3a182e9b56a61b163d28b7f2 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 14 Aug 2026 21:28:21 +0000 Subject: [PATCH 12/12] other subset of changes --- cpp/benchmarks/io/parquet/experimental/variant/extract.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp index 784346f88b3d..184822687cae 100644 --- a/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp +++ b/cpp/benchmarks/io/parquet/experimental/variant/extract.cpp @@ -62,8 +62,9 @@ constexpr uint8_t make_variant_primitive_header(variant_primitive_type type) } // Header byte for a short string of the given length (must fit in 6 bits: 0..63). -constexpr uint8_t make_variant_short_string_header(std::size_t length) +uint8_t make_variant_short_string_header(std::size_t length) { + CUDF_EXPECTS(length <= 63, "Short string length must fit in 6 bits (0..63)"); return make_variant_header(variant_basic_type::SHORT_STRING, static_cast(length)); } @@ -79,6 +80,7 @@ constexpr uint8_t make_variant_array_header() return make_variant_header(variant_basic_type::ARRAY, 0); } +// Append the low `width` bytes of `bits` to `out` in little-endian order. void append_le(std::vector& out, uint64_t bits, int width) { for (int i = 0; i < width; ++i) {