From ca6beca7068b79976b72592bed6ea9c4d6cdbeda Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 22 Jul 2026 17:01:16 +0000 Subject: [PATCH 01/14] first draft --- cpp/CMakeLists.txt | 1 + cpp/include/cudf/io/experimental/variant.hpp | 51 ++ .../io/parquet/experimental/variant_encode.cu | 559 ++++++++++++++++++ cpp/tests/CMakeLists.txt | 1 + .../io/experimental/variant_encode_test.cpp | 556 +++++++++++++++++ 5 files changed, 1168 insertions(+) create mode 100644 cpp/src/io/parquet/experimental/variant_encode.cu create mode 100644 cpp/tests/io/experimental/variant_encode_test.cpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index f542bc4b6dc8..4147a373343a 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -796,6 +796,7 @@ add_library( src/io/parquet/experimental/hybrid_scan_preprocess.cu src/io/parquet/experimental/page_index_filter.cu src/io/parquet/experimental/page_index_filter_utils.cu + src/io/parquet/experimental/variant_encode.cu src/io/parquet/experimental/variant_extract.cu src/io/parquet/experimental/variant_path.cpp src/io/parquet/expression_transform_helpers.cpp diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 4ee7747dc055..7d931b30a4bf 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -7,13 +7,16 @@ #include #include +#include #include #include #include +#include #include #include +#include #include /** @@ -107,6 +110,54 @@ namespace io::parquet::experimental { rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); +/** + * @brief Encode a STRING column as a scalar VARIANT column. + * + * Each non-null string is encoded as a VARIANT scalar value: a short_string blob (1 header byte + + * payload) for strings shorter than 64 bytes, or a long_string primitive blob (1 header + 4-byte + * LE length + payload) for longer strings. The metadata blob for every row is the minimal + * empty-dictionary encoding `{0x01, 0x00, 0x00}` (version 1, 0 keys). Null input rows produce + * null VARIANT struct rows. + * + * @param strings STRING column to encode + * @param stream CUDA stream + * @param mr Device memory resource + * @return VARIANT column: `STRUCT, list>` (metadata child, value child) + * + * @throws std::invalid_argument if `strings` is not a STRING column + */ +[[nodiscard]] std::unique_ptr encode_strings_to_variant( + column_view const& strings, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Encode a table as a VARIANT object column, one row per VARIANT. + * + * Each table row is encoded as a VARIANT object with one field per column. The shared metadata + * blob stores column names as a sorted UTF-8 key dictionary (version 1, offset_size chosen to + * fit the total key-string length). Null column values are encoded as VARIANT null primitives. + * + * Supported column types: `INT8`, `INT16`, `INT32`, `INT64`, `STRING`, and columns whose type is + * `EMPTY` (treated as all-null). Other types throw. Tables must have fewer than 256 columns. + * Total encoded output must be < 2 GiB. + * + * @param input Table to encode (typically produced by `cudf::io::read_json`) + * @param column_names Column name for each column in `input`; must satisfy + * `column_names.size() == input.num_columns()` + * @param stream CUDA stream + * @param mr Device memory resource + * @return VARIANT column: `STRUCT, list>` (metadata child, value child) + * + * @throws std::invalid_argument if any column has an unsupported type, if the table has ≥ 256 + * columns, or if `column_names.size() != input.num_columns()` + */ +[[nodiscard]] std::unique_ptr encode_variant( + cudf::table_view const& input, + cudf::host_span column_names, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + /** @} */ } // namespace io::parquet::experimental } // namespace CUDF_EXPORT cudf diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu new file mode 100644 index 000000000000..66bcc4082b87 --- /dev/null +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -0,0 +1,559 @@ +/* + * 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 + +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include +#include + +namespace cudf { +namespace io::parquet::experimental { +namespace { + +// ────────────────────────────────────────────────────────────────────────────── +// Encoding constants (from the VARIANT spec; mirrors variant_extract.cu) +// ────────────────────────────────────────────────────────────────────────────── + +constexpr uint8_t k_null_value = 0x00; +constexpr uint8_t k_int8_header = 0x0c; +constexpr uint8_t k_int16_header = 0x10; +constexpr uint8_t k_int32_header = 0x14; +constexpr uint8_t k_int64_header = 0x18; +constexpr uint8_t k_long_string_header = 0x40; // primitive type 16 + +// short_string: basic_type=1, length in bits 2..7 of the value_metadata byte +constexpr uint8_t k_short_string_basic_type = 0x01; + +// Object with field_id_size=1, field_offset_size=4, num_elements_size=1: +// value_header = (is_large=0 << 4) | (field_id_size-1=0 << 2) | (field_offset_size-1=3) = 0x03 +// value_metadata = (value_header << 2) | basic_type::object(2) = 0x0e +constexpr uint8_t k_object_value_metadata = 0x0e; + +// Bytes in the constant empty-dictionary metadata blob {version=1, 0 keys} +constexpr size_type k_empty_meta_size = 3; + +constexpr int block_size = 256; + +// ────────────────────────────────────────────────────────────────────────────── +// Per-field device helpers +// ────────────────────────────────────────────────────────────────────────────── + +__device__ size_type field_encoded_size(column_device_view const& col, size_type row) +{ + if (col.type().id() == type_id::EMPTY || !col.is_valid(row)) { return 1; } + switch (col.type().id()) { + case type_id::INT8: return 2; + case type_id::INT16: return 3; + case type_id::INT32: return 5; + case type_id::INT64: return 9; + case type_id::STRING: { + auto const len = static_cast(col.element(row).size_bytes()); + return len < 64 ? 1 + len : 5 + len; + } + default: return 1; + } +} + +// Write the encoded bytes for one field into `out`, returns bytes written. +__device__ size_type write_field_value(uint8_t* out, column_device_view const& col, size_type row) +{ + if (col.type().id() == type_id::EMPTY || !col.is_valid(row)) { + out[0] = k_null_value; + return 1; + } + switch (col.type().id()) { + case type_id::INT8: { + auto const v = col.element(row); + out[0] = k_int8_header; + cuda::std::memcpy(out + 1, &v, 1); + return 2; + } + case type_id::INT16: { + auto const v = col.element(row); + out[0] = k_int16_header; + cuda::std::memcpy(out + 1, &v, 2); + return 3; + } + case type_id::INT32: { + auto const v = col.element(row); + out[0] = k_int32_header; + cuda::std::memcpy(out + 1, &v, 4); + return 5; + } + case type_id::INT64: { + auto const v = col.element(row); + out[0] = k_int64_header; + cuda::std::memcpy(out + 1, &v, 8); + return 9; + } + case type_id::STRING: { + auto const sv = col.element(row); + auto const len = static_cast(sv.size_bytes()); + if (len < 64) { + out[0] = static_cast(k_short_string_basic_type | (len << 2)); + cuda::std::memcpy(out + 1, sv.data(), len); + return 1 + len; + } else { + out[0] = k_long_string_header; + uint32_t const u32 = static_cast(len); + cuda::std::memcpy(out + 1, &u32, 4); + cuda::std::memcpy(out + 5, sv.data(), len); + return 5 + len; + } + } + default: out[0] = k_null_value; return 1; + } +} + +// ────────────────────────────────────────────────────────────────────────────── +// encode_strings_to_variant kernels +// ────────────────────────────────────────────────────────────────────────────── + +// Pass 1: compute per-row value blob sizes (metadata is always k_empty_meta_size = 3 bytes). +CUDF_KERNEL __launch_bounds__(block_size) void string_value_sizes_kernel( + column_device_view strings, device_span d_val_sizes) +{ + auto const num_rows = strings.size(); + auto const tid = cudf::detail::grid_1d::global_thread_id(); + auto const stride = cudf::detail::grid_1d::grid_stride(); + + for (auto row = tid; row < num_rows; row += stride) { + if (!strings.is_valid(row)) { + d_val_sizes[row] = 1; + continue; + } + auto const len = static_cast(strings.element(row).size_bytes()); + d_val_sizes[row] = len < 64 ? 1 + len : 5 + len; + } +} + +// Pass 2: write metadata and value blobs. +// Metadata is constant {0x01, 0x00, 0x00} at stride k_empty_meta_size bytes per row. +CUDF_KERNEL __launch_bounds__(block_size) void string_encode_write_kernel( + column_device_view strings, + device_span d_val_offsets, + uint8_t* d_val_buf, + uint8_t* d_meta_buf) +{ + auto const num_rows = strings.size(); + auto const tid = cudf::detail::grid_1d::global_thread_id(); + auto const stride = cudf::detail::grid_1d::grid_stride(); + + for (auto row = tid; row < num_rows; row += stride) { + // Constant metadata: {version=1, num_keys=0, sentinel_offset=0} + auto* mp = d_meta_buf + row * k_empty_meta_size; + mp[0] = 0x01; + mp[1] = 0x00; + mp[2] = 0x00; + + // Value + uint8_t* vp = d_val_buf + d_val_offsets[row]; + if (!strings.is_valid(row)) { + vp[0] = k_null_value; + continue; + } + auto const sv = strings.element(row); + auto const len = static_cast(sv.size_bytes()); + if (len < 64) { + vp[0] = static_cast(k_short_string_basic_type | (len << 2)); + cuda::std::memcpy(vp + 1, sv.data(), len); + } else { + vp[0] = k_long_string_header; + uint32_t const u32 = static_cast(len); + cuda::std::memcpy(vp + 1, &u32, 4); + cuda::std::memcpy(vp + 5, sv.data(), len); + } + } +} + +// ────────────────────────────────────────────────────────────────────────────── +// encode_variant kernels +// ────────────────────────────────────────────────────────────────────────────── + +// Pass 1: compute per-row value blob sizes. +CUDF_KERNEL __launch_bounds__(block_size) void object_value_sizes_kernel( + cudf::table_device_view tbl, + device_span sort_order, + device_span d_val_sizes) +{ + auto const num_rows = static_cast(d_val_sizes.size()); + auto const N = tbl.num_columns(); + auto const tid = cudf::detail::grid_1d::global_thread_id(); + auto const stride = cudf::detail::grid_1d::grid_stride(); + + for (auto row = tid; row < num_rows; row += stride) { + // Header: value_metadata(1) + num_elements(1) + field_ids(N) + field_offsets((N+1)*4) + size_type size = 2 + N + (N + 1) * 4; + for (int i = 0; i < N; i++) { + size += field_encoded_size(tbl.column(sort_order[i]), row); + } + d_val_sizes[row] = size; + } +} + +// Pass 2: write metadata and value blobs. +// Metadata is constant, stored at stride `meta_size` bytes per row. +CUDF_KERNEL __launch_bounds__(block_size) void object_encode_write_kernel( + cudf::table_device_view tbl, + device_span sort_order, + device_span d_val_offsets, + uint8_t* d_val_buf, + uint8_t* d_meta_buf, + uint8_t const* d_meta_template, + size_type meta_size) +{ + auto const num_rows = static_cast(d_val_offsets.size()) - 1; + auto const N = tbl.num_columns(); + auto const tid = cudf::detail::grid_1d::global_thread_id(); + auto const stride = cudf::detail::grid_1d::grid_stride(); + + for (auto row = tid; row < num_rows; row += stride) { + // Broadcast constant metadata template + cuda::std::memcpy(d_meta_buf + row * meta_size, d_meta_template, meta_size); + + // Write the VARIANT object value + uint8_t* p = d_val_buf + d_val_offsets[row]; + + *p++ = k_object_value_metadata; + *p++ = static_cast(N); + + // field_ids: sorted column IDs are 0..N-1 (dictionary is in sorted order) + for (int i = 0; i < N; i++) { + *p++ = static_cast(i); + } + + // Compute per-field sizes (stack array, safe for N < 256) + uint32_t field_sizes[256]; + for (int i = 0; i < N; i++) { + field_sizes[i] = static_cast(field_encoded_size(tbl.column(sort_order[i]), row)); + } + + // Write (N+1) field offsets, 4 bytes each (LE) + uint32_t running = 0; + for (int i = 0; i <= N; i++) { + cuda::std::memcpy(p, &running, 4); + p += 4; + if (i < N) { running += field_sizes[i]; } + } + + // Write field values in sorted order + for (int i = 0; i < N; i++) { + p += write_field_value(p, tbl.column(sort_order[i]), row); + } + } +} + +// ────────────────────────────────────────────────────────────────────────────── +// Host: build metadata blob and column sort order +// ────────────────────────────────────────────────────────────────────────────── + +// Returns (metadata_bytes, sort_order) where sort_order[i] is the original column index of +// the i-th lexicographically sorted column. Column names appear in sorted order in the blob. +std::pair, std::vector> build_metadata_blob( + cudf::host_span column_names) +{ + auto const N = static_cast(column_names.size()); + + std::vector sort_order(N); + std::iota(sort_order.begin(), sort_order.end(), 0); + std::stable_sort(sort_order.begin(), sort_order.end(), [&](int a, int b) { + return column_names[a] < column_names[b]; + }); + + // Choose offset_size based on total UTF-8 key length + std::size_t total_key_bytes = 0; + for (auto const& name : column_names) { + total_key_bytes += name.size(); + } + int const offset_size = total_key_bytes <= 0xFFu ? 1 : total_key_bytes <= 0xFFFFu ? 2 : 4; + + // header: bits[7:6] = offset_size-1, bits[3:0] = version=1 + uint8_t const header = static_cast(((offset_size - 1) << 6) | 0x01u); + + std::vector blob; + blob.push_back(header); + + // num_keys (offset_size bytes LE) + for (int b = 0; b < offset_size; b++) { + blob.push_back(static_cast((static_cast(N) >> (8 * b)) & 0xFFu)); + } + + // offsets[0..N] (offset_size bytes each LE), relative to start of string_data + std::size_t running = 0; + auto push_offset = [&](std::size_t v) { + for (int b = 0; b < offset_size; b++) { + blob.push_back(static_cast((v >> (8 * b)) & 0xFFu)); + } + }; + push_offset(0); + for (int i = 0; i < N; i++) { + running += column_names[sort_order[i]].size(); + push_offset(running); + } + + // string_data: keys in sorted order + for (int i = 0; i < N; i++) { + auto const& name = column_names[sort_order[i]]; + blob.insert(blob.end(), name.begin(), name.end()); + } + + return {std::move(blob), std::move(sort_order)}; +} + +// Build a list column for constant-stride blobs (all rows the same M bytes). +// Offsets are [0, M, 2M, ..., num_rows*M]; data is filled by caller via kernel. +std::pair, std::unique_ptr> make_constant_list_buffers( + size_type num_rows, size_type M, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +{ + auto offsets = make_numeric_column( + data_type{type_id::INT32}, num_rows + 1, mask_state::UNALLOCATED, stream, mr); + thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + offsets->mutable_view().begin(), + offsets->mutable_view().end(), + int32_t{0}, + static_cast(M)); + + auto data = make_numeric_column( + data_type{type_id::UINT8}, num_rows * M, mask_state::UNALLOCATED, stream, mr); + return {std::move(offsets), std::move(data)}; +} + +} // namespace + +namespace detail { + +// ────────────────────────────────────────────────────────────────────────────── +// encode_strings_to_variant +// ────────────────────────────────────────────────────────────────────────────── + +std::unique_ptr encode_strings_to_variant(column_view const& strings, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(strings.type().id() == type_id::STRING, + "encode_strings_to_variant: input must be a STRING column", + std::invalid_argument); + + auto const num_rows = strings.size(); + + auto make_empty_list = [&] { + return make_lists_column( + 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); + }; + + if (num_rows == 0) { + std::vector> ch; + ch.push_back(make_empty_list()); + ch.push_back(make_empty_list()); + return make_structs_column(0, std::move(ch), 0, {}, stream, mr); + } + + auto strings_dv = column_device_view::create(strings, stream); + + // ── Metadata: constant k_empty_meta_size bytes per row ── + auto [meta_offsets_col, meta_data_col] = + make_constant_list_buffers(num_rows, k_empty_meta_size, stream, mr); + + // ── Value sizes ── + rmm::device_uvector d_val_sizes(num_rows, stream, mr); + { + cudf::detail::grid_1d grid{num_rows, block_size}; + string_value_sizes_kernel<<>>(*strings_dv, + d_val_sizes); + CUDF_CUDA_TRY(cudaGetLastError()); + } + + auto [val_offsets_col, total_val_bytes] = cudf::strings::detail::make_offsets_child_column( + d_val_sizes.begin(), d_val_sizes.end(), stream, mr); + CUDF_EXPECTS(total_val_bytes <= std::numeric_limits::max(), + "VARIANT value bytes exceed 2 GiB limit", + std::overflow_error); + + auto val_data_col = make_numeric_column(data_type{type_id::UINT8}, + static_cast(total_val_bytes), + mask_state::UNALLOCATED, + stream, + mr); + + // ── Write pass (metadata + values) ── + { + device_span d_val_offsets{val_offsets_col->view().data(), + static_cast(num_rows + 1)}; + cudf::detail::grid_1d grid{num_rows, block_size}; + string_encode_write_kernel<<>>( + *strings_dv, + d_val_offsets, + val_data_col->mutable_view().data(), + meta_data_col->mutable_view().data()); + CUDF_CUDA_TRY(cudaGetLastError()); + } + + // ── Null mask: propagate from input ── + size_type const null_count = strings.null_count(); + rmm::device_buffer null_mask = + null_count > 0 ? cudf::detail::copy_bitmask(strings, stream, mr) : rmm::device_buffer{}; + + // ── Assemble STRUCT, list> ── + auto meta_col = + make_lists_column(num_rows, std::move(meta_offsets_col), std::move(meta_data_col), 0, {}); + auto val_col = + make_lists_column(num_rows, std::move(val_offsets_col), std::move(val_data_col), 0, {}); + + std::vector> children; + children.push_back(std::move(meta_col)); + children.push_back(std::move(val_col)); + return make_structs_column( + num_rows, std::move(children), null_count, std::move(null_mask), stream, mr); +} + +// ────────────────────────────────────────────────────────────────────────────── +// encode_variant +// ────────────────────────────────────────────────────────────────────────────── + +std::unique_ptr encode_variant(cudf::table_view const& input, + cudf::host_span column_names, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto const N = input.num_columns(); + auto const num_rows = input.num_rows(); + + CUDF_EXPECTS(static_cast(N) == column_names.size(), + "encode_variant: column_names.size() must equal input.num_columns()", + std::invalid_argument); + CUDF_EXPECTS( + N < 256, "encode_variant: table must have fewer than 256 columns", std::invalid_argument); + + for (int i = 0; i < N; i++) { + auto const id = input.column(i).type().id(); + CUDF_EXPECTS( + id == type_id::EMPTY || id == type_id::INT8 || id == type_id::INT16 || id == type_id::INT32 || + id == type_id::INT64 || id == type_id::STRING, + "encode_variant: unsupported column type — supported: EMPTY, INT8/16/32/64, STRING", + std::invalid_argument); + } + + auto make_empty_list = [&] { + return make_lists_column( + 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); + }; + + if (num_rows == 0) { + std::vector> ch; + ch.push_back(make_empty_list()); + ch.push_back(make_empty_list()); + return make_structs_column(0, std::move(ch), 0, {}, stream, mr); + } + + // ── Metadata blob and sort order ── + auto [meta_bytes, sort_order] = build_metadata_blob(column_names); + auto const M = static_cast(meta_bytes.size()); + + rmm::device_uvector d_meta_template = + cudf::detail::make_device_uvector_async(meta_bytes, stream, mr); + + rmm::device_uvector d_sort_order = + cudf::detail::make_device_uvector_async(sort_order, stream, mr); + + // ── Column device views on device (via table_device_view) ── + auto d_table = cudf::table_device_view::create(input, stream); + + // ── Metadata buffers: constant M bytes per row ── + auto [meta_offsets_col, meta_data_col] = make_constant_list_buffers(num_rows, M, stream, mr); + + // ── Value sizes (pass 1) ── + rmm::device_uvector d_val_sizes(num_rows, stream, mr); + { + cudf::detail::grid_1d grid{num_rows, block_size}; + object_value_sizes_kernel<<>>( + *d_table, d_sort_order, d_val_sizes); + CUDF_CUDA_TRY(cudaGetLastError()); + } + + auto [val_offsets_col, total_val_bytes] = cudf::strings::detail::make_offsets_child_column( + d_val_sizes.begin(), d_val_sizes.end(), stream, mr); + CUDF_EXPECTS(total_val_bytes <= std::numeric_limits::max(), + "VARIANT value bytes exceed 2 GiB limit", + std::overflow_error); + + auto val_data_col = make_numeric_column(data_type{type_id::UINT8}, + static_cast(total_val_bytes), + mask_state::UNALLOCATED, + stream, + mr); + + // ── Write pass (pass 2): metadata + values ── + { + device_span d_val_offsets{val_offsets_col->view().data(), + static_cast(num_rows + 1)}; + cudf::detail::grid_1d grid{num_rows, block_size}; + object_encode_write_kernel<<>>( + *d_table, + d_sort_order, + d_val_offsets, + val_data_col->mutable_view().data(), + meta_data_col->mutable_view().data(), + d_meta_template.data(), + M); + CUDF_CUDA_TRY(cudaGetLastError()); + } + + // ── Assemble STRUCT, list> (output rows are never null) ── + auto meta_col = + make_lists_column(num_rows, std::move(meta_offsets_col), std::move(meta_data_col), 0, {}); + auto val_col = + make_lists_column(num_rows, std::move(val_offsets_col), std::move(val_data_col), 0, {}); + + std::vector> children; + children.push_back(std::move(meta_col)); + children.push_back(std::move(val_col)); + return make_structs_column(num_rows, std::move(children), 0, {}, stream, mr); +} + +} // namespace detail + +std::unique_ptr encode_strings_to_variant(column_view const& strings, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::encode_strings_to_variant(strings, stream, mr); +} + +std::unique_ptr encode_variant(cudf::table_view const& input, + cudf::host_span column_names, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::encode_variant(input, column_names, stream, mr); +} + +} // namespace io::parquet::experimental +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 25938c6b5e05..61674ef6b8df 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -359,6 +359,7 @@ ConfigureTest( io/parquet_common.cpp io/parquet_test.cpp ) +ConfigureTest(VARIANT_ENCODE_TEST io/experimental/variant_encode_test.cpp) ConfigureTest(VARIANT_EXTRACT_TEST io/experimental/variant_extract_test.cpp) ConfigureTest( PARQUET_DELETION_VECTORS_TEST diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp new file mode 100644 index 000000000000..fc9ddc200b3e --- /dev/null +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -0,0 +1,556 @@ +/* + * 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 + +// ────────────────────────────────────────────────────────────────────────────── +// Test helpers +// ────────────────────────────────────────────────────────────────────────────── + +namespace { + +// Byte-level encoding helpers (host-side, matching the GPU implementation) + +inline std::vector enc_null() { return {0x00}; } + +inline std::vector enc_int8(int8_t v) { return {0x0c, static_cast(v)}; } + +inline std::vector enc_int16(int16_t v) +{ + auto const u = static_cast(v); + return {0x10, static_cast(u & 0xffu), static_cast((u >> 8) & 0xffu)}; +} + +inline std::vector enc_int32(int32_t v) +{ + auto const u = static_cast(v); + return {0x14, + static_cast(u & 0xffu), + static_cast((u >> 8) & 0xffu), + static_cast((u >> 16) & 0xffu), + static_cast((u >> 24) & 0xffu)}; +} + +inline std::vector enc_int64(int64_t v) +{ + auto const u = static_cast(v); + std::vector out{0x18}; + for (int b = 0; b < 8; b++) { + out.push_back(static_cast((u >> (8 * b)) & 0xffu)); + } + return out; +} + +inline std::vector enc_short_string(std::string_view s) +{ + std::vector out{static_cast(0x01 | (s.size() << 2))}; + out.insert(out.end(), s.begin(), s.end()); + return out; +} + +inline std::vector enc_long_string(std::string_view s) +{ + auto const len = static_cast(s.size()); + std::vector out{0x40}; + for (int b = 0; b < 4; b++) { + out.push_back(static_cast((len >> (8 * b)) & 0xffu)); + } + out.insert(out.end(), s.begin(), s.end()); + return out; +} + +inline std::vector enc_string(std::string_view s) +{ + return s.size() < 64 ? enc_short_string(s) : enc_long_string(s); +} + +// Build VARIANT metadata blob for a sorted list of key names (offset_size=1 assumed; keys must +// have total byte length <= 255). +inline std::vector build_metadata(std::vector const& sorted_keys) +{ + std::vector out{0x01, static_cast(sorted_keys.size())}; + std::vector offsets{0x00}; + uint8_t running = 0; + for (auto const& k : sorted_keys) { + running = static_cast(running + k.size()); + offsets.push_back(running); + } + out.insert(out.end(), offsets.begin(), offsets.end()); + for (auto const& k : sorted_keys) { + out.insert(out.end(), k.begin(), k.end()); + } + return out; +} + +// Build an object VARIANT value blob with field_id_size=1, field_offset_size=4. +// `field_values` must be in the same order as sorted_keys (IDs 0..N-1). +inline std::vector build_object_value( + std::vector> const& field_values) +{ + auto const N = static_cast(field_values.size()); + std::vector out; + + // value_metadata = 0x0e: object, field_id_size=1, field_offset_size=4, num_elements_size=1 + out.push_back(0x0e); + out.push_back(static_cast(N)); + + // field_ids = 0, 1, ..., N-1 + for (int i = 0; i < N; i++) { + out.push_back(static_cast(i)); + } + + // field_offsets[0..N] (4 bytes each LE) + uint32_t running = 0; + for (int i = 0; i <= N; i++) { + for (int b = 0; b < 4; b++) { + out.push_back(static_cast((running >> (8 * b)) & 0xffu)); + } + if (i < N) { running += static_cast(field_values[i].size()); } + } + + // field values + for (auto const& fv : field_values) { + out.insert(out.end(), fv.begin(), fv.end()); + } + return out; +} + +// Build a multi-row VARIANT struct column from explicit per-row (meta, value) byte vectors. +inline cudf::test::structs_column_wrapper make_variant_column( + std::vector> const& meta_rows, + std::vector> const& val_rows) +{ + auto build = [](std::vector> const& rows) { + auto n = static_cast(rows.size()); + std::vector offsets(n + 1, 0); + std::vector flat; + for (int i = 0; i < n; i++) { + flat.insert(flat.end(), rows[i].begin(), rows[i].end()); + offsets[i + 1] = static_cast(flat.size()); + } + auto offs_col = + cudf::test::fixed_width_column_wrapper(offsets.begin(), offsets.end()).release(); + auto data_col = + cudf::test::fixed_width_column_wrapper(flat.begin(), flat.end()).release(); + return cudf::make_lists_column(n, std::move(offs_col), std::move(data_col), 0, {}); + }; + + std::vector> children; + children.push_back(build(meta_rows)); + children.push_back(build(val_rows)); + return cudf::test::structs_column_wrapper{std::move(children)}; +} + +// Return the metadata child column view (child 0) of a VARIANT struct column. +inline cudf::column_view meta_child(cudf::column_view const& variant_col) +{ + return cudf::structs_column_view{variant_col}.child(0); +} + +// Return the value child column view (child 1) of a VARIANT struct column. +inline cudf::column_view val_child(cudf::column_view const& variant_col) +{ + return cudf::structs_column_view{variant_col}.child(1); +} + +} // namespace + +// ────────────────────────────────────────────────────────────────────────────── +// encode_strings_to_variant tests +// ────────────────────────────────────────────────────────────────────────────── + +struct EncodeStringsToVariantTest : public cudf::test::BaseFixture {}; + +TEST_F(EncodeStringsToVariantTest, EmptyInput) +{ + cudf::test::strings_column_wrapper input{}; + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + EXPECT_EQ(got->type().id(), cudf::type_id::STRUCT); + EXPECT_EQ(got->size(), 0); + EXPECT_EQ(got->null_count(), 0); +} + +TEST_F(EncodeStringsToVariantTest, WrongTypeThrows) +{ + cudf::test::fixed_width_column_wrapper input{1, 2, 3}; + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream())), + std::invalid_argument); +} + +TEST_F(EncodeStringsToVariantTest, ShortString) +{ + cudf::test::strings_column_wrapper input{"hi"}; + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 0); + + auto const expected_val = enc_short_string("hi"); + auto const expected_meta = std::vector{0x01, 0x00, 0x00}; + auto expected = make_variant_column({expected_meta}, {expected_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeStringsToVariantTest, LongString) +{ + // A string of exactly 64 bytes should use long_string encoding + std::string const s(64, 'x'); + cudf::test::strings_column_wrapper input{s}; + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + auto const expected_val = enc_long_string(s); + auto const expected_meta = std::vector{0x01, 0x00, 0x00}; + auto expected = make_variant_column({expected_meta}, {expected_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeStringsToVariantTest, NullInputProducesNullStructRow) +{ + cudf::test::strings_column_wrapper input({"hello", "", "world"}, {true, false, true}); + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 3); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(EncodeStringsToVariantTest, AllNullInput) +{ + cudf::test::strings_column_wrapper input({"", "", ""}, {false, false, false}); + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + EXPECT_EQ(got->null_count(), 3); +} + +TEST_F(EncodeStringsToVariantTest, MultiRowMixedLengths) +{ + std::string const short_s = "abc"; // len 3 < 64 + std::string const long_s = std::string(100, 'z'); // len 100 >= 64 + + cudf::test::strings_column_wrapper input{short_s, long_s}; + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 2); + EXPECT_EQ(got->null_count(), 0); + + // Verify the value child has the right encoding for each row + auto const ev0 = enc_short_string(short_s); + auto const ev1 = enc_long_string(long_s); + auto expected = make_variant_column( + {std::vector{0x01, 0x00, 0x00}, std::vector{0x01, 0x00, 0x00}}, {ev0, ev1}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeStringsToVariantTest, RoundtripWithCastVariant) +{ + // encode then decode: cast_variant should recover the original strings + cudf::test::strings_column_wrapper input{"foo", "bar", "baz"}; + auto variant = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::cast_variant( + val_child(*variant), cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, input); +} + +TEST_F(EncodeStringsToVariantTest, RoundtripLongStringWithCastVariant) +{ + std::string const long_s = std::string(128, 'a'); + cudf::test::strings_column_wrapper input{long_s}; + auto variant = cudf::io::parquet::experimental::encode_strings_to_variant( + input, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::cast_variant( + val_child(*variant), cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, input); +} + +// ────────────────────────────────────────────────────────────────────────────── +// encode_variant tests +// ────────────────────────────────────────────────────────────────────────────── + +struct EncodeVariantTest : public cudf::test::BaseFixture {}; + +TEST_F(EncodeVariantTest, EmptyTable) +{ + cudf::test::fixed_width_column_wrapper col{}; + cudf::table_view tbl{{col}}; + std::vector names{"x"}; + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->type().id(), cudf::type_id::STRUCT); + EXPECT_EQ(got->size(), 0); +} + +TEST_F(EncodeVariantTest, UnsupportedTypeThrows) +{ + cudf::test::fixed_width_column_wrapper col{1.0f, 2.0f}; + cudf::table_view tbl{{col}}; + std::vector names{"f"}; + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_variant( + tbl, names, cudf::test::get_default_stream())), + std::invalid_argument); +} + +TEST_F(EncodeVariantTest, ColumnNamesMismatchThrows) +{ + cudf::test::fixed_width_column_wrapper col{1, 2}; + cudf::table_view tbl{{col}}; + std::vector names{"a", "b"}; // 2 names, 1 column + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_variant( + tbl, names, cudf::test::get_default_stream())), + std::invalid_argument); +} + +TEST_F(EncodeVariantTest, SingleInt32Column) +{ + // Table: col "x" = [42] + cudf::test::fixed_width_column_wrapper col{42}; + cudf::table_view tbl{{col}}; + std::vector names{"x"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 0); + + // metadata: {0x01, num_keys=1, offsets=[0,1], "x"} = {0x01, 0x01, 0x00, 0x01, 'x'} + auto const exp_meta = build_metadata({"x"}); + // value: object with 1 field, field_id=0, field_value=enc_int32(42) + auto const exp_val = build_object_value({enc_int32(42)}); + + auto expected = make_variant_column({exp_meta}, {exp_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, SingleStringColumn) +{ + cudf::test::strings_column_wrapper col{"hello"}; + cudf::table_view tbl{{col}}; + std::vector names{"s"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto const exp_meta = build_metadata({"s"}); + auto const exp_val = build_object_value({enc_short_string("hello")}); + + auto expected = make_variant_column({exp_meta}, {exp_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, NullValueEncodedAsVariantNull) +{ + // A null INT32 value should produce a null primitive (0x00) inside the object + cudf::test::fixed_width_column_wrapper col({0}, {false}); + cudf::table_view tbl{{col}}; + std::vector names{"x"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->null_count(), 0); // struct row is never null + + auto const exp_meta = build_metadata({"x"}); + auto const exp_val = build_object_value({enc_null()}); + + auto expected = make_variant_column({exp_meta}, {exp_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, MultipleColumns_SortedNames) +{ + // Columns "b", "a": the output dictionary must sort to ["a","b"] + cudf::test::fixed_width_column_wrapper col_b{10}; + cudf::test::strings_column_wrapper col_a{"hi"}; + cudf::table_view tbl{{col_b, col_a}}; + std::vector names{"b", "a"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + // Dictionary is sorted: ["a", "b"]. sort_order = [1 (col "a"), 0 (col "b")]. + // Field 0 (id=0, name="a") → col_a value = enc_short_string("hi") + // Field 1 (id=1, name="b") → col_b value = enc_int32(10) + auto const exp_meta = build_metadata({"a", "b"}); + auto const exp_val = build_object_value({enc_short_string("hi"), enc_int32(10)}); + + auto expected = make_variant_column({exp_meta}, {exp_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, AllIntegerTypes) +{ + cudf::test::fixed_width_column_wrapper c8{int8_t{7}}; + cudf::test::fixed_width_column_wrapper c16{int16_t{300}}; + cudf::test::fixed_width_column_wrapper c32{int32_t{70000}}; + cudf::test::fixed_width_column_wrapper c64{int64_t{5000000000LL}}; + cudf::table_view tbl{{c8, c16, c32, c64}}; + std::vector names{"i8", "i16", "i32", "i64"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 0); + + // sorted names: i16, i32, i64, i8 + auto const exp_meta = build_metadata({"i16", "i32", "i64", "i8"}); + auto const exp_val = + build_object_value({enc_int16(300), enc_int32(70000), enc_int64(5000000000LL), enc_int8(7)}); + + auto expected = make_variant_column({exp_meta}, {exp_val}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, MultiRow) +{ + // 3 rows: col "a" (int32), col "b" (string) + cudf::test::fixed_width_column_wrapper col_a{1, 2, 3}; + cudf::test::strings_column_wrapper col_b{"x", "yy", "zzz"}; + cudf::table_view tbl{{col_a, col_b}}; + std::vector names{"a", "b"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 3); + EXPECT_EQ(got->null_count(), 0); + + auto const meta = build_metadata({"a", "b"}); + auto const exp_val0 = build_object_value({enc_int32(1), enc_short_string("x")}); + auto const exp_val1 = build_object_value({enc_int32(2), enc_short_string("yy")}); + auto const exp_val2 = build_object_value({enc_int32(3), enc_short_string("zzz")}); + + auto expected = make_variant_column({meta, meta, meta}, {exp_val0, exp_val1, exp_val2}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, MixedNullsInMultiColumn) +{ + // col "a": [1, null, 3] + // col "b": ["x", "y", null] + cudf::test::fixed_width_column_wrapper col_a({1, 0, 3}, {true, false, true}); + cudf::test::strings_column_wrapper col_b({"x", "y", ""}, {true, true, false}); + cudf::table_view tbl{{col_a, col_b}}; + std::vector names{"a", "b"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->null_count(), 0); // struct rows never null in encode_variant + + auto const meta = build_metadata({"a", "b"}); + auto const exp_val0 = build_object_value({enc_int32(1), enc_short_string("x")}); + auto const exp_val1 = build_object_value({enc_null(), enc_short_string("y")}); + auto const exp_val2 = build_object_value({enc_int32(3), enc_null()}); + + auto expected = make_variant_column({meta, meta, meta}, {exp_val0, exp_val1, exp_val2}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, RoundtripInt32WithExtractVariant) +{ + // Encode, then extract "x" and decode as INT32 + cudf::test::fixed_width_column_wrapper col{10, 20, 30}; + cudf::table_view tbl{{col}}; + std::vector names{"x"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *variant, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} + +TEST_F(EncodeVariantTest, RoundtripStringWithExtractVariant) +{ + cudf::test::strings_column_wrapper col{"alpha", "beta", "gamma"}; + cudf::table_view tbl{{col}}; + std::vector names{"s"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *variant, "s", cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} + +TEST_F(EncodeVariantTest, RoundtripNullsWithExtractVariant) +{ + cudf::test::fixed_width_column_wrapper col({5, 0, 15}, {true, false, true}); + cudf::table_view tbl{{col}}; + std::vector names{"v"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *variant, "v", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} + +TEST_F(EncodeVariantTest, RoundtripMultiColumnWithExtractVariant) +{ + cudf::test::fixed_width_column_wrapper col_a{100, 200, 300}; + cudf::test::strings_column_wrapper col_b{"p", "q", "r"}; + cudf::table_view tbl{{col_a, col_b}}; + std::vector names{"a", "b"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded_a = cudf::io::parquet::experimental::extract_variant_field( + *variant, "a", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto decoded_b = cudf::io::parquet::experimental::extract_variant_field( + *variant, "b", cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded_a, col_a); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded_b, col_b); +} + +TEST_F(EncodeVariantTest, LargeStringFieldRoundtrip) +{ + // Strings longer than 64 bytes must use long_string encoding in the object field + std::string const long_s(200, 'k'); + cudf::test::strings_column_wrapper col{long_s}; + cudf::table_view tbl{{col}}; + std::vector names{"big"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *variant, "big", cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} From 692a4776a09905e9b90825e41363b94dedaf3832 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 21:57:53 +0000 Subject: [PATCH 02/14] changes --- cpp/include/cudf/io/experimental/variant.hpp | 5 +- .../io/parquet/experimental/variant_encode.cu | 50 +++++---- .../io/experimental/variant_encode_test.cpp | 24 ++--- .../io/experimental/variant_extract_test.cpp | 102 ++++++++++++++++++ 4 files changed, 148 insertions(+), 33 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 68bfe21913be..d292256da719 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -7,6 +7,7 @@ #include #include +#include #include #include #include @@ -121,7 +122,7 @@ namespace io::parquet::experimental { * empty-dictionary encoding `{0x01, 0x00, 0x00}` (version 1, 0 keys). Null input rows produce * null VARIANT struct rows. * - * @param strings STRING column to encode + * @param input STRING column to encode * @param stream CUDA stream * @param mr Device memory resource * @return VARIANT column: `STRUCT, list>` (metadata child, value child) @@ -129,7 +130,7 @@ namespace io::parquet::experimental { * @throws std::invalid_argument if `strings` is not a STRING column */ [[nodiscard]] std::unique_ptr encode_strings_to_variant( - column_view const& strings, + cudf::strings_column_view const& input, rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index 66bcc4082b87..a8385623272e 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -45,6 +46,8 @@ constexpr uint8_t k_int8_header = 0x0c; constexpr uint8_t k_int16_header = 0x10; constexpr uint8_t k_int32_header = 0x14; constexpr uint8_t k_int64_header = 0x18; +constexpr uint8_t k_float64_header = 0x1c; // primitive type 7 +constexpr uint8_t k_float32_header = 0x38; // primitive type 14 constexpr uint8_t k_long_string_header = 0x40; // primitive type 16 // short_string: basic_type=1, length in bits 2..7 of the value_metadata byte @@ -72,6 +75,8 @@ __device__ size_type field_encoded_size(column_device_view const& col, size_type case type_id::INT16: return 3; case type_id::INT32: return 5; case type_id::INT64: return 9; + case type_id::FLOAT32: return 5; + case type_id::FLOAT64: return 9; case type_id::STRING: { auto const len = static_cast(col.element(row).size_bytes()); return len < 64 ? 1 + len : 5 + len; @@ -112,6 +117,18 @@ __device__ size_type write_field_value(uint8_t* out, column_device_view const& c cuda::std::memcpy(out + 1, &v, 8); return 9; } + case type_id::FLOAT32: { + auto const v = col.element(row); + out[0] = k_float32_header; + cuda::std::memcpy(out + 1, &v, 4); + return 5; + } + case type_id::FLOAT64: { + auto const v = col.element(row); + out[0] = k_float64_header; + cuda::std::memcpy(out + 1, &v, 8); + return 9; + } case type_id::STRING: { auto const sv = col.element(row); auto const len = static_cast(sv.size_bytes()); @@ -352,15 +369,11 @@ namespace detail { // encode_strings_to_variant // ────────────────────────────────────────────────────────────────────────────── -std::unique_ptr encode_strings_to_variant(column_view const& strings, +std::unique_ptr encode_strings_to_variant(cudf::strings_column_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - CUDF_EXPECTS(strings.type().id() == type_id::STRING, - "encode_strings_to_variant: input must be a STRING column", - std::invalid_argument); - - auto const num_rows = strings.size(); + auto const num_rows = input.size(); auto make_empty_list = [&] { return make_lists_column( @@ -374,7 +387,7 @@ std::unique_ptr encode_strings_to_variant(column_view const& strings, return make_structs_column(0, std::move(ch), 0, {}, stream, mr); } - auto strings_dv = column_device_view::create(strings, stream); + auto input_dv = column_device_view::create(input.parent(), stream); // ── Metadata: constant k_empty_meta_size bytes per row ── auto [meta_offsets_col, meta_data_col] = @@ -384,7 +397,7 @@ std::unique_ptr encode_strings_to_variant(column_view const& strings, rmm::device_uvector d_val_sizes(num_rows, stream, mr); { cudf::detail::grid_1d grid{num_rows, block_size}; - string_value_sizes_kernel<<>>(*strings_dv, + string_value_sizes_kernel<<>>(*input_dv, d_val_sizes); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -407,7 +420,7 @@ std::unique_ptr encode_strings_to_variant(column_view const& strings, static_cast(num_rows + 1)}; cudf::detail::grid_1d grid{num_rows, block_size}; string_encode_write_kernel<<>>( - *strings_dv, + *input_dv, d_val_offsets, val_data_col->mutable_view().data(), meta_data_col->mutable_view().data()); @@ -415,9 +428,9 @@ std::unique_ptr encode_strings_to_variant(column_view const& strings, } // ── Null mask: propagate from input ── - size_type const null_count = strings.null_count(); + size_type const null_count = input.null_count(); rmm::device_buffer null_mask = - null_count > 0 ? cudf::detail::copy_bitmask(strings, stream, mr) : rmm::device_buffer{}; + null_count > 0 ? cudf::detail::copy_bitmask(input.parent(), stream, mr) : rmm::device_buffer{}; // ── Assemble STRUCT, list> ── auto meta_col = @@ -452,11 +465,12 @@ std::unique_ptr encode_variant(cudf::table_view const& input, for (int i = 0; i < N; i++) { auto const id = input.column(i).type().id(); - CUDF_EXPECTS( - id == type_id::EMPTY || id == type_id::INT8 || id == type_id::INT16 || id == type_id::INT32 || - id == type_id::INT64 || id == type_id::STRING, - "encode_variant: unsupported column type — supported: EMPTY, INT8/16/32/64, STRING", - std::invalid_argument); + CUDF_EXPECTS(id == type_id::EMPTY || id == type_id::INT8 || id == type_id::INT16 || + id == type_id::INT32 || id == type_id::INT64 || id == type_id::FLOAT32 || + id == type_id::FLOAT64 || id == type_id::STRING, + "encode_variant: unsupported column type — supported: EMPTY, INT8/16/32/64, " + "FLOAT32/64, STRING", + std::invalid_argument); } auto make_empty_list = [&] { @@ -538,12 +552,12 @@ std::unique_ptr encode_variant(cudf::table_view const& input, } // namespace detail -std::unique_ptr encode_strings_to_variant(column_view const& strings, +std::unique_ptr encode_strings_to_variant(cudf::strings_column_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); - return detail::encode_strings_to_variant(strings, stream, mr); + return detail::encode_strings_to_variant(input, stream, mr); } std::unique_ptr encode_variant(cudf::table_view const& input, diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index fc9ddc200b3e..1354e19478b4 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -179,7 +179,7 @@ TEST_F(EncodeStringsToVariantTest, EmptyInput) { cudf::test::strings_column_wrapper input{}; auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); EXPECT_EQ(got->type().id(), cudf::type_id::STRUCT); EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); @@ -188,16 +188,14 @@ TEST_F(EncodeStringsToVariantTest, EmptyInput) TEST_F(EncodeStringsToVariantTest, WrongTypeThrows) { cudf::test::fixed_width_column_wrapper input{1, 2, 3}; - EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream())), - std::invalid_argument); + EXPECT_THROW(static_cast(cudf::strings_column_view{input}), cudf::logic_error); } TEST_F(EncodeStringsToVariantTest, ShortString) { cudf::test::strings_column_wrapper input{"hi"}; auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); EXPECT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 0); @@ -214,7 +212,7 @@ TEST_F(EncodeStringsToVariantTest, LongString) std::string const s(64, 'x'); cudf::test::strings_column_wrapper input{s}; auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); auto const expected_val = enc_long_string(s); auto const expected_meta = std::vector{0x01, 0x00, 0x00}; @@ -226,7 +224,7 @@ TEST_F(EncodeStringsToVariantTest, NullInputProducesNullStructRow) { cudf::test::strings_column_wrapper input({"hello", "", "world"}, {true, false, true}); auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); EXPECT_EQ(got->size(), 3); EXPECT_EQ(got->null_count(), 1); @@ -236,7 +234,7 @@ TEST_F(EncodeStringsToVariantTest, AllNullInput) { cudf::test::strings_column_wrapper input({"", "", ""}, {false, false, false}); auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); EXPECT_EQ(got->null_count(), 3); } @@ -248,7 +246,7 @@ TEST_F(EncodeStringsToVariantTest, MultiRowMixedLengths) cudf::test::strings_column_wrapper input{short_s, long_s}; auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); EXPECT_EQ(got->size(), 2); EXPECT_EQ(got->null_count(), 0); @@ -266,7 +264,7 @@ TEST_F(EncodeStringsToVariantTest, RoundtripWithCastVariant) // encode then decode: cast_variant should recover the original strings cudf::test::strings_column_wrapper input{"foo", "bar", "baz"}; auto variant = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); auto decoded = cudf::io::parquet::experimental::cast_variant( val_child(*variant), cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); @@ -279,7 +277,7 @@ TEST_F(EncodeStringsToVariantTest, RoundtripLongStringWithCastVariant) std::string const long_s = std::string(128, 'a'); cudf::test::strings_column_wrapper input{long_s}; auto variant = cudf::io::parquet::experimental::encode_strings_to_variant( - input, cudf::test::get_default_stream()); + cudf::strings_column_view{input}, cudf::test::get_default_stream()); auto decoded = cudf::io::parquet::experimental::cast_variant( val_child(*variant), cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); @@ -307,9 +305,9 @@ TEST_F(EncodeVariantTest, EmptyTable) TEST_F(EncodeVariantTest, UnsupportedTypeThrows) { - cudf::test::fixed_width_column_wrapper col{1.0f, 2.0f}; + cudf::test::fixed_width_column_wrapper col{1, 2}; cudf::table_view tbl{{col}}; - std::vector names{"f"}; + std::vector names{"u"}; EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_variant( tbl, names, cudf::test::get_default_stream())), std::invalid_argument); diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 4356fb40d1c9..d6c118105f26 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -883,3 +883,105 @@ TEST_F(CastVariantTest, EmptyInput) EXPECT_EQ(got->null_count(), 0); } } + +struct EncodeVariantTest : public cudf::test::BaseFixture {}; + +TEST_F(EncodeVariantTest, Float32RoundTrip) +{ + auto const stream = cudf::test::get_default_stream(); + + cudf::test::fixed_width_column_wrapper floats({1.5f, -2.25f, 3.14159f}); + cudf::table_view tbl{{floats}}; + std::vector names{"f"}; + + auto encoded = cudf::io::parquet::experimental::encode_variant(tbl, names, stream); + auto got = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "f", cudf::data_type{cudf::type_id::FLOAT32}, stream); + + cudf::test::fixed_width_column_wrapper expected({1.5f, -2.25f, 3.14159f}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, Float64RoundTrip) +{ + auto const stream = cudf::test::get_default_stream(); + + cudf::test::fixed_width_column_wrapper doubles({1.5, -2.25, 3.141592653589793}); + cudf::table_view tbl{{doubles}}; + std::vector names{"d"}; + + auto encoded = cudf::io::parquet::experimental::encode_variant(tbl, names, stream); + auto got = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "d", cudf::data_type{cudf::type_id::FLOAT64}, stream); + + cudf::test::fixed_width_column_wrapper expected({1.5, -2.25, 3.141592653589793}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, FloatNullsRoundTrip) +{ + auto const stream = cudf::test::get_default_stream(); + + cudf::test::fixed_width_column_wrapper floats({1.0f, 0.0f, -3.5f}, {true, false, true}); + cudf::table_view tbl{{floats}}; + std::vector names{"f"}; + + auto encoded = cudf::io::parquet::experimental::encode_variant(tbl, names, stream); + auto got = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "f", cudf::data_type{cudf::type_id::FLOAT32}, stream); + + cudf::test::fixed_width_column_wrapper expected({1.0f, 0.0f, -3.5f}, {true, false, true}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, MixedFloatAndIntColumns) +{ + auto const stream = cudf::test::get_default_stream(); + + cudf::test::fixed_width_column_wrapper ints({10, 20, 30}); + cudf::test::fixed_width_column_wrapper floats({1.5f, 2.5f, 3.5f}); + cudf::test::fixed_width_column_wrapper doubles({10.1, 20.2, 30.3}); + cudf::table_view tbl{{ints, floats, doubles}}; + std::vector names{"i", "f", "d"}; + + auto encoded = cudf::io::parquet::experimental::encode_variant(tbl, names, stream); + + auto got_i = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "i", cudf::data_type{cudf::type_id::INT32}, stream); + auto got_f = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "f", cudf::data_type{cudf::type_id::FLOAT32}, stream); + auto got_d = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "d", cudf::data_type{cudf::type_id::FLOAT64}, stream); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got_i, + cudf::test::fixed_width_column_wrapper({10, 20, 30})); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got_f, + cudf::test::fixed_width_column_wrapper({1.5f, 2.5f, 3.5f})); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + *got_d, cudf::test::fixed_width_column_wrapper({10.1, 20.2, 30.3})); +} + +TEST_F(EncodeVariantTest, FloatEncodeMatchesApacheBytes) +{ + // Verify that encode_variant produces the exact bytes from the Apache parquet-testing fixtures + // for float32 (1234567936.0f) and float64 (1234567890.1234). + auto const stream = cudf::test::get_default_stream(); + + cudf::test::fixed_width_column_wrapper f32_col({1234567936.0f}); + cudf::test::fixed_width_column_wrapper f64_col({1234567890.1234}); + cudf::table_view tbl{{f32_col, f64_col}}; + std::vector names{"float_field", "double_field"}; + + auto encoded = cudf::io::parquet::experimental::encode_variant(tbl, names, stream); + + // Extract and cast back — the decoded values must match the originals exactly. + auto got_f32 = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "float_field", cudf::data_type{cudf::type_id::FLOAT32}, stream); + auto got_f64 = cudf::io::parquet::experimental::extract_variant_field( + *encoded, "double_field", cudf::data_type{cudf::type_id::FLOAT64}, stream); + + cudf::test::fixed_width_column_wrapper expected_f32({1234567936.0f}); + cudf::test::fixed_width_column_wrapper expected_f64({1234567890.1234}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got_f32, expected_f32); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got_f64, expected_f64); +} From 1ced02f6a09507b1d6dddf285a94c849a900477a Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 17:56:36 -0500 Subject: [PATCH 03/14] Update cpp/include/cudf/io/experimental/variant.hpp Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com> --- cpp/include/cudf/io/experimental/variant.hpp | 21 ++++++++++---------- 1 file changed, 11 insertions(+), 10 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index d292256da719..2dca8225dce1 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -135,25 +135,26 @@ namespace io::parquet::experimental { rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** - * @brief Encode a table as a VARIANT object column, one row per VARIANT. + * `@brief` Encode a table as a VARIANT object column, one row per VARIANT. * * Each table row is encoded as a VARIANT object with one field per column. The shared metadata * blob stores column names as a sorted UTF-8 key dictionary (version 1, offset_size chosen to * fit the total key-string length). Null column values are encoded as VARIANT null primitives. * - * Supported column types: `INT8`, `INT16`, `INT32`, `INT64`, `STRING`, and columns whose type is - * `EMPTY` (treated as all-null). Other types throw. Tables must have fewer than 256 columns. - * Total encoded output must be < 2 GiB. + * Supported column types: `INT8`, `INT16`, `INT32`, `INT64`, `FLOAT32`, `FLOAT64`, `STRING`, and + * columns whose type is `EMPTY` (treated as all-null). Other types throw. Tables must have fewer + * than 256 columns. * - * @param input Table to encode (typically produced by `cudf::io::read_json`) - * @param column_names Column name for each column in `input`; must satisfy + * `@param` input Table to encode (typically produced by `cudf::io::read_json`) + * `@param` column_names Column name for each column in `input`; must satisfy * `column_names.size() == input.num_columns()` - * @param stream CUDA stream - * @param mr Device memory resource - * @return VARIANT column: `STRUCT, list>` (metadata child, value child) + * `@param` stream CUDA stream + * `@param` mr Device memory resource + * `@return` VARIANT column: `STRUCT, list>` (metadata child, value child) * - * @throws std::invalid_argument if any column has an unsupported type, if the table has ≥ 256 + * `@throws` std::invalid_argument if any column has an unsupported type, if the table has ≥ 256 * columns, or if `column_names.size() != input.num_columns()` + * `@throws` std::overflow_error if the total encoded value bytes exceed 2 GiB */ [[nodiscard]] std::unique_ptr encode_variant( cudf::table_view const& input, From 8edfbbb59d1b2be3744ced4ff5c529a92331dc7e Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 17:57:00 -0500 Subject: [PATCH 04/14] Update cpp/src/io/parquet/experimental/variant_encode.cu Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com> --- cpp/src/io/parquet/experimental/variant_encode.cu | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index a8385623272e..a4c8879398b7 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -376,8 +376,13 @@ std::unique_ptr encode_strings_to_variant(cudf::strings_column_view cons auto const num_rows = input.size(); auto make_empty_list = [&] { - return make_lists_column( - 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); + return make_lists_column(0, + make_empty_column(type_id::INT32, stream, mr), + make_empty_column(type_id::UINT8, stream, mr), + 0, + {}, + stream, + mr); }; if (num_rows == 0) { From 00611b73ba033c199bb40c348053e9e679b158b3 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 17:57:11 -0500 Subject: [PATCH 05/14] Update cpp/tests/io/experimental/variant_encode_test.cpp Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com> --- .../io/experimental/variant_encode_test.cpp | 27 ++++++++++--------- 1 file changed, 14 insertions(+), 13 deletions(-) diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index 1354e19478b4..6b8c9f7cfd15 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -3,19 +3,20 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include -#include -#include -#include - -#include -#include -#include -#include - -#include -#include -#include +`#include` +`#include` +`#include` +`#include` +`#include` + +`#include` +`#include` +`#include` +`#include` + +`#include` +`#include` +`#include` // ────────────────────────────────────────────────────────────────────────────── // Test helpers From c2c7086a0b6f6f8102b148a0e5da45f405542380 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 23:13:19 +0000 Subject: [PATCH 06/14] fixing ai comments & merge --- cpp/include/cudf/io/experimental/variant.hpp | 2 +- .../io/parquet/experimental/variant_encode.cu | 17 +++++++++++------ .../io/experimental/variant_extract_test.cpp | 6 +++--- 3 files changed, 15 insertions(+), 10 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index d292256da719..0765b45e5af5 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -127,7 +127,7 @@ namespace io::parquet::experimental { * @param mr Device memory resource * @return VARIANT column: `STRUCT, list>` (metadata child, value child) * - * @throws std::invalid_argument if `strings` is not a STRING column + * @throws cudf::logic_error if `input` is not a STRING column */ [[nodiscard]] std::unique_ptr encode_strings_to_variant( cudf::strings_column_view const& input, diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index a8385623272e..9f513a4f9d22 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -348,6 +348,9 @@ std::pair, std::vector> build_metadata_blob( std::pair, std::unique_ptr> make_constant_list_buffers( size_type num_rows, size_type M, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { + CUDF_EXPECTS(static_cast(num_rows) * M <= std::numeric_limits::max(), + "metadata size exceeds 2 GiB limit", + std::overflow_error); auto offsets = make_numeric_column( data_type{type_id::INT32}, num_rows + 1, mask_state::UNALLOCATED, stream, mr); thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -394,7 +397,8 @@ std::unique_ptr encode_strings_to_variant(cudf::strings_column_view cons make_constant_list_buffers(num_rows, k_empty_meta_size, stream, mr); // ── Value sizes ── - rmm::device_uvector d_val_sizes(num_rows, stream, mr); + rmm::device_uvector d_val_sizes( + num_rows, stream, cudf::get_current_device_resource_ref()); { cudf::detail::grid_1d grid{num_rows, block_size}; string_value_sizes_kernel<<>>(*input_dv, @@ -489,11 +493,11 @@ std::unique_ptr encode_variant(cudf::table_view const& input, auto [meta_bytes, sort_order] = build_metadata_blob(column_names); auto const M = static_cast(meta_bytes.size()); - rmm::device_uvector d_meta_template = - cudf::detail::make_device_uvector_async(meta_bytes, stream, mr); + rmm::device_uvector d_meta_template = cudf::detail::make_device_uvector_async( + meta_bytes, stream, cudf::get_current_device_resource_ref()); - rmm::device_uvector d_sort_order = - cudf::detail::make_device_uvector_async(sort_order, stream, mr); + rmm::device_uvector d_sort_order = cudf::detail::make_device_uvector_async( + sort_order, stream, cudf::get_current_device_resource_ref()); // ── Column device views on device (via table_device_view) ── auto d_table = cudf::table_device_view::create(input, stream); @@ -502,7 +506,8 @@ std::unique_ptr encode_variant(cudf::table_view const& input, auto [meta_offsets_col, meta_data_col] = make_constant_list_buffers(num_rows, M, stream, mr); // ── Value sizes (pass 1) ── - rmm::device_uvector d_val_sizes(num_rows, stream, mr); + rmm::device_uvector d_val_sizes( + num_rows, stream, cudf::get_current_device_resource_ref()); { cudf::detail::grid_1d grid{num_rows, block_size}; object_value_sizes_kernel<<>>( diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index b13e31c0dfd7..8d9235a6f687 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1379,10 +1379,10 @@ TEST_F(EncodeVariantTest, MixedFloatAndIntColumns) *got_d, cudf::test::fixed_width_column_wrapper({10.1, 20.2, 30.3})); } -TEST_F(EncodeVariantTest, FloatEncodeMatchesApacheBytes) +TEST_F(EncodeVariantTest, ApacheFixtureValuesRoundTrip) { - // Verify that encode_variant produces the exact bytes from the Apache parquet-testing fixtures - // for float32 (1234567936.0f) and float64 (1234567890.1234). + // Round-trip the specific float32/float64 values from the Apache parquet-testing fixtures + // (1234567936.0f and 1234567890.1234) through encode_variant + extract_variant_field. auto const stream = cudf::test::get_default_stream(); cudf::test::fixed_width_column_wrapper f32_col({1234567936.0f}); From df45af3d06224754d04f346e659db8ba612d2316 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 23:21:20 +0000 Subject: [PATCH 07/14] build --- .../io/parquet/experimental/variant_encode.cu | 9 ++---- .../io/experimental/variant_encode_test.cpp | 28 +++++++++---------- 2 files changed, 16 insertions(+), 21 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index 44425a2a9f60..9f513a4f9d22 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -379,13 +379,8 @@ std::unique_ptr encode_strings_to_variant(cudf::strings_column_view cons auto const num_rows = input.size(); auto make_empty_list = [&] { - return make_lists_column(0, - make_empty_column(type_id::INT32, stream, mr), - make_empty_column(type_id::UINT8, stream, mr), - 0, - {}, - stream, - mr); + return make_lists_column( + 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); }; if (num_rows == 0) { diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index 6b8c9f7cfd15..7d7b4d95dbff 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -3,20 +3,20 @@ * SPDX-License-Identifier: Apache-2.0 */ -`#include` -`#include` -`#include` -`#include` -`#include` - -`#include` -`#include` -`#include` -`#include` - -`#include` -`#include` -`#include` +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include // ────────────────────────────────────────────────────────────────────────────── // Test helpers From 8089b45ffe6d74aa204dbb9d07e77f8caeaa8396 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 23:30:11 +0000 Subject: [PATCH 08/14] changers --- .../io/experimental/variant_extract_test.cpp | 42 +++++++++++++++---- 1 file changed, 35 insertions(+), 7 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 8d9235a6f687..f01b5f406a93 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -443,6 +443,13 @@ inline std::vector enc_int64(int64_t v) return out; } +inline std::vector enc_float32(float v) +{ + std::vector out{make_variant_primitive(variant_primitive_type::FLOAT32)}; + append_le(out, std::bit_cast(v), 4); + return out; +} + inline std::vector enc_float64(double v) { std::vector out{make_variant_primitive(variant_primitive_type::FLOAT64)}; @@ -1034,9 +1041,9 @@ TEST_F(CastVariantTest, EmptyInput) TEST_F(CastVariantTest, CastToUnsupportedTargetThrows) { - // cast_variant only supports INT8/16/32/64 and STRING targets. Every other target is rejected at - // compile-time dispatch on the requested output type, independent of the input bytes, so a single - // well-formed placeholder row triggers the same throw for all of them. + // cast_variant supports INT8/16/32/64, FLOAT32/64, and STRING targets. Every other target is + // rejected at compile-time dispatch on the requested output type, independent of the input bytes, + // so a single well-formed placeholder row triggers the same throw for all of them. auto stream = cudf::test::get_default_stream(); std::vector const val{make_variant_primitive(variant_primitive_type::NULLVAL)}; cudf::test::lists_column_wrapper values(val.begin(), val.end()); @@ -1064,10 +1071,11 @@ TEST_F(CastVariantTest, CastToUnsupportedTargetThrows) TEST_F(CastVariantTest, CastSourceTargetMatrix) { // Exhaustively covers (source physical type) x (supported target) casts. The supported targets - // are INT8/16/32/64 and STRING. Expected behaviour: - // - integer targets: only a source whose physical type has the *exact* same width decodes; - // every - // other source (including narrower/wider ints) yields null — cast_variant does not widen. + // are INT8/16/32/64, FLOAT32/64, and STRING. Expected behaviour: + // - integer/float targets: only a source whose physical type has the *exact* same width + // decodes; + // every other source (including narrower/wider types) yields null — cast_variant does not + // widen or convert between ints and floats. // - STRING target: short_string and long_string sources decode; every other source yields null. auto const stream = cudf::test::get_default_stream(); @@ -1083,6 +1091,7 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) {"int16", enc_int16(1234)}, {"int32", enc_int32(123456)}, {"int64", enc_int64(1234567890123456789LL)}, + {"float32", enc_float32(1.5f)}, {"float64", enc_float64(2.5)}, {"short_string", enc_short_string("hi")}, {"long_string", enc_long_string(std::string(70, 'a'))}, @@ -1113,6 +1122,25 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) check_int_target.template operator()("int32", int32_t{123456}); check_int_target.template operator()("int64", int64_t{1234567890123456789LL}); + // Float targets: exact-width match only; no int<->float conversion. + auto check_float_target = [&](char const* match_label, T match_value) { + auto const target = cudf::data_type{cudf::type_to_id()}; + for (auto const& src : sources) { + SCOPED_TRACE(std::string{"float target "} + match_label + ", source " + src.label); + auto values = values_of(src.bytes); + auto got = cudf::io::parquet::experimental::cast_variant(values, target, stream); + if (std::string_view{src.label} == match_label) { + cudf::test::fixed_width_column_wrapper const expected{match_value}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } else { + ASSERT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 1); + } + } + }; + check_float_target.template operator()("float32", float{1.5f}); + check_float_target.template operator()("float64", double{2.5}); + // STRING target: short_string and long_string decode; every other source is null. auto const string_type = cudf::data_type{cudf::type_id::STRING}; for (auto const& src : sources) { From 66df6ad04da7e1b046e153514b6e2d286ccd802d Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 30 Jul 2026 22:56:05 +0000 Subject: [PATCH 09/14] reviews --- .../io/parquet/experimental/variant_encode.cu | 6 +++++ .../io/experimental/variant_encode_test.cpp | 24 +++++++++++++++++++ 2 files changed, 30 insertions(+) diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index 9f513a4f9d22..a9dbdb14062a 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -303,6 +303,12 @@ std::pair, std::vector> build_metadata_blob( return column_names[a] < column_names[b]; }); + for (int i = 1; i < N; i++) { + CUDF_EXPECTS(column_names[sort_order[i]] != column_names[sort_order[i - 1]], + "encode_variant: duplicate column names are not allowed", + std::invalid_argument); + } + // Choose offset_size based on total UTF-8 key length std::size_t total_key_bytes = 0; for (auto const& name : column_names) { diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index 7d7b4d95dbff..bcf7df65a7db 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -537,6 +537,30 @@ TEST_F(EncodeVariantTest, RoundtripMultiColumnWithExtractVariant) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded_b, col_b); } +TEST_F(EncodeVariantTest, DuplicateColumnNamesThrows) +{ + cudf::test::fixed_width_column_wrapper col_a{1}; + cudf::test::fixed_width_column_wrapper col_b{2}; + cudf::table_view tbl{{col_a, col_b}}; + std::vector names{"a", "a"}; + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_variant( + tbl, names, cudf::test::get_default_stream())), + std::invalid_argument); +} + +TEST_F(EncodeVariantTest, DuplicateColumnNamesAfterSortThrows) +{ + // Names that sort into adjacent positions: ["b", "a", "a"] → sorted ["a", "a", "b"] + cudf::test::fixed_width_column_wrapper col_b{1}; + cudf::test::fixed_width_column_wrapper col_a0{2}; + cudf::test::fixed_width_column_wrapper col_a1{3}; + cudf::table_view tbl{{col_b, col_a0, col_a1}}; + std::vector names{"b", "a", "a"}; + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::encode_variant( + tbl, names, cudf::test::get_default_stream())), + std::invalid_argument); +} + TEST_F(EncodeVariantTest, LargeStringFieldRoundtrip) { // Strings longer than 64 bytes must use long_string encoding in the object field From f551764dbfdb1f8f8ca460ac1cbd0b202e521d14 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 30 Jul 2026 23:13:18 +0000 Subject: [PATCH 10/14] coderabbit --- .../io/parquet/experimental/variant_encode.cu | 19 +++++++++++++------ 1 file changed, 13 insertions(+), 6 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index a9dbdb14062a..6c2ed7719aa0 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -303,12 +303,6 @@ std::pair, std::vector> build_metadata_blob( return column_names[a] < column_names[b]; }); - for (int i = 1; i < N; i++) { - CUDF_EXPECTS(column_names[sort_order[i]] != column_names[sort_order[i - 1]], - "encode_variant: duplicate column names are not allowed", - std::invalid_argument); - } - // Choose offset_size based on total UTF-8 key length std::size_t total_key_bytes = 0; for (auto const& name : column_names) { @@ -483,6 +477,19 @@ std::unique_ptr encode_variant(cudf::table_view const& input, std::invalid_argument); } + { + std::vector name_order(N); + std::iota(name_order.begin(), name_order.end(), 0); + std::stable_sort(name_order.begin(), name_order.end(), [&](int a, int b) { + return column_names[a] < column_names[b]; + }); + for (int i = 1; i < N; i++) { + CUDF_EXPECTS(column_names[name_order[i]] != column_names[name_order[i - 1]], + "encode_variant: duplicate column names are not allowed", + std::invalid_argument); + } + } + auto make_empty_list = [&] { return make_lists_column( 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); From f50b7b3a9d756bd3efbb5390bdf30d6f85478acb Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 4 Aug 2026 23:17:38 +0000 Subject: [PATCH 11/14] addressing reviewers --- cpp/include/cudf/io/experimental/variant.hpp | 16 ++--- .../io/parquet/experimental/variant_encode.cu | 61 ++++++++--------- .../io/experimental/variant_encode_test.cpp | 67 +++++++++++++++++-- 3 files changed, 100 insertions(+), 44 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index f15059b5001d..378831dde749 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -135,7 +135,7 @@ namespace io::parquet::experimental { rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** - * `@brief` Encode a table as a VARIANT object column, one row per VARIANT. + * @brief Encode a table as a VARIANT object column, one row per VARIANT. * * Each table row is encoded as a VARIANT object with one field per column. The shared metadata * blob stores column names as a sorted UTF-8 key dictionary (version 1, offset_size chosen to @@ -145,16 +145,16 @@ namespace io::parquet::experimental { * columns whose type is `EMPTY` (treated as all-null). Other types throw. Tables must have fewer * than 256 columns. * - * `@param` input Table to encode (typically produced by `cudf::io::read_json`) - * `@param` column_names Column name for each column in `input`; must satisfy + * @param input Table to encode (typically produced by `cudf::io::read_json`) + * @param column_names Column name for each column in `input`; must satisfy * `column_names.size() == input.num_columns()` - * `@param` stream CUDA stream - * `@param` mr Device memory resource - * `@return` VARIANT column: `STRUCT, list>` (metadata child, value child) + * @param stream CUDA stream + * @param mr Device memory resource + * @return VARIANT column: `STRUCT, list>` (metadata child, value child) * - * `@throws` std::invalid_argument if any column has an unsupported type, if the table has ≥ 256 + * @throws std::invalid_argument if any column has an unsupported type, if the table has ≥ 256 * columns, or if `column_names.size() != input.num_columns()` - * `@throws` std::overflow_error if the total encoded value bytes exceed 2 GiB + * @throws std::overflow_error if the total encoded value bytes exceed 2 GiB */ [[nodiscard]] std::unique_ptr encode_variant( cudf::table_view const& input, diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index 6c2ed7719aa0..13e1ade5def1 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -67,7 +67,7 @@ constexpr int block_size = 256; // Per-field device helpers // ────────────────────────────────────────────────────────────────────────────── -__device__ size_type field_encoded_size(column_device_view const& col, size_type row) +__device__ int64_t field_encoded_size(column_device_view const& col, size_type row) { if (col.type().id() == type_id::EMPTY || !col.is_valid(row)) { return 1; } switch (col.type().id()) { @@ -78,7 +78,7 @@ __device__ size_type field_encoded_size(column_device_view const& col, size_type case type_id::FLOAT32: return 5; case type_id::FLOAT64: return 9; case type_id::STRING: { - auto const len = static_cast(col.element(row).size_bytes()); + auto const len = static_cast(col.element(row).size_bytes()); return len < 64 ? 1 + len : 5 + len; } default: return 1; @@ -165,8 +165,12 @@ CUDF_KERNEL __launch_bounds__(block_size) void string_value_sizes_kernel( d_val_sizes[row] = 1; continue; } - auto const len = static_cast(strings.element(row).size_bytes()); - d_val_sizes[row] = len < 64 ? 1 + len : 5 + len; + auto const len = static_cast(strings.element(row).size_bytes()); + auto const enc_sz = len < 64 ? 1 + len : 5 + len; + d_val_sizes[row] = static_cast( + enc_sz <= static_cast(cuda::std::numeric_limits::max()) + ? enc_sz + : cuda::std::numeric_limits::max()); } } @@ -226,11 +230,14 @@ CUDF_KERNEL __launch_bounds__(block_size) void object_value_sizes_kernel( for (auto row = tid; row < num_rows; row += stride) { // Header: value_metadata(1) + num_elements(1) + field_ids(N) + field_offsets((N+1)*4) - size_type size = 2 + N + (N + 1) * 4; + int64_t size = static_cast(2 + N + (N + 1) * 4); for (int i = 0; i < N; i++) { size += field_encoded_size(tbl.column(sort_order[i]), row); } - d_val_sizes[row] = size; + // Saturate to INT32_MAX so the host-side total-bytes overflow check fires. + d_val_sizes[row] = size <= static_cast(cuda::std::numeric_limits::max()) + ? static_cast(size) + : cuda::std::numeric_limits::max(); } } @@ -265,18 +272,16 @@ CUDF_KERNEL __launch_bounds__(block_size) void object_encode_write_kernel( *p++ = static_cast(i); } - // Compute per-field sizes (stack array, safe for N < 256) - uint32_t field_sizes[256]; - for (int i = 0; i < N; i++) { - field_sizes[i] = static_cast(field_encoded_size(tbl.column(sort_order[i]), row)); - } - - // Write (N+1) field offsets, 4 bytes each (LE) + // Write (N+1) field offsets, 4 bytes each (LE). + // field_encoded_size is called once per field here (and again below in write_field_value), + // which avoids the 1 KiB field_sizes[256] stack array that would spill to local memory. uint32_t running = 0; for (int i = 0; i <= N; i++) { cuda::std::memcpy(p, &running, 4); p += 4; - if (i < N) { running += field_sizes[i]; } + if (i < N) { + running += static_cast(field_encoded_size(tbl.column(sort_order[i]), row)); + } } // Write field values in sorted order @@ -408,7 +413,10 @@ std::unique_ptr encode_strings_to_variant(cudf::strings_column_view cons auto [val_offsets_col, total_val_bytes] = cudf::strings::detail::make_offsets_child_column( d_val_sizes.begin(), d_val_sizes.end(), stream, mr); - CUDF_EXPECTS(total_val_bytes <= std::numeric_limits::max(), + // INT64 offsets are produced when total_bytes >= get_offset64_threshold() (== INT32_MAX). + // Checking the column type is more robust than comparing total_val_bytes to INT32_MAX, + // since the boundary case (total == INT32_MAX) produces INT64 but the numeric check passes. + CUDF_EXPECTS(val_offsets_col->type().id() == type_id::INT32, "VARIANT value bytes exceed 2 GiB limit", std::overflow_error); @@ -477,17 +485,12 @@ std::unique_ptr encode_variant(cudf::table_view const& input, std::invalid_argument); } - { - std::vector name_order(N); - std::iota(name_order.begin(), name_order.end(), 0); - std::stable_sort(name_order.begin(), name_order.end(), [&](int a, int b) { - return column_names[a] < column_names[b]; - }); - for (int i = 1; i < N; i++) { - CUDF_EXPECTS(column_names[name_order[i]] != column_names[name_order[i - 1]], - "encode_variant: duplicate column names are not allowed", - std::invalid_argument); - } + // ── Metadata blob, sort order, and duplicate check (sort once) ── + auto [meta_bytes, sort_order] = build_metadata_blob(column_names); + for (int i = 1; i < N; i++) { + CUDF_EXPECTS(column_names[sort_order[i]] != column_names[sort_order[i - 1]], + "encode_variant: duplicate column names are not allowed", + std::invalid_argument); } auto make_empty_list = [&] { @@ -502,9 +505,7 @@ std::unique_ptr encode_variant(cudf::table_view const& input, return make_structs_column(0, std::move(ch), 0, {}, stream, mr); } - // ── Metadata blob and sort order ── - auto [meta_bytes, sort_order] = build_metadata_blob(column_names); - auto const M = static_cast(meta_bytes.size()); + auto const M = static_cast(meta_bytes.size()); rmm::device_uvector d_meta_template = cudf::detail::make_device_uvector_async( meta_bytes, stream, cudf::get_current_device_resource_ref()); @@ -530,7 +531,7 @@ std::unique_ptr encode_variant(cudf::table_view const& input, auto [val_offsets_col, total_val_bytes] = cudf::strings::detail::make_offsets_child_column( d_val_sizes.begin(), d_val_sizes.end(), stream, mr); - CUDF_EXPECTS(total_val_bytes <= std::numeric_limits::max(), + CUDF_EXPECTS(val_offsets_col->type().id() == type_id::INT32, "VARIANT value bytes exceed 2 GiB limit", std::overflow_error); diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index bcf7df65a7db..6f78ed073423 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -10,11 +10,13 @@ #include #include +#include #include #include #include #include +#include #include #include @@ -186,12 +188,6 @@ TEST_F(EncodeStringsToVariantTest, EmptyInput) EXPECT_EQ(got->null_count(), 0); } -TEST_F(EncodeStringsToVariantTest, WrongTypeThrows) -{ - cudf::test::fixed_width_column_wrapper input{1, 2, 3}; - EXPECT_THROW(static_cast(cudf::strings_column_view{input}), cudf::logic_error); -} - TEST_F(EncodeStringsToVariantTest, ShortString) { cudf::test::strings_column_wrapper input{"hi"}; @@ -286,6 +282,25 @@ TEST_F(EncodeStringsToVariantTest, RoundtripLongStringWithCastVariant) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, input); } +TEST_F(EncodeStringsToVariantTest, SlicedInput) +{ + // Rows 0 and 3 are padding; the slice covers rows [1, 3) to exercise a non-zero column offset. + cudf::test::strings_column_wrapper full{"ignore", "hello", "world", "ignore"}; + auto const sliced = cudf::slice(full, {1, 3})[0]; + + auto got = cudf::io::parquet::experimental::encode_strings_to_variant( + cudf::strings_column_view{sliced}, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 2); + EXPECT_EQ(got->null_count(), 0); + + auto const meta = std::vector{0x01, 0x00, 0x00}; + auto const exp_v0 = enc_short_string("hello"); + auto const exp_v1 = enc_short_string("world"); + auto expected = make_variant_column({meta, meta}, {exp_v0, exp_v1}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + // ────────────────────────────────────────────────────────────────────────────── // encode_variant tests // ────────────────────────────────────────────────────────────────────────────── @@ -577,3 +592,43 @@ TEST_F(EncodeVariantTest, LargeStringFieldRoundtrip) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); } + +TEST_F(EncodeVariantTest, MultiBlock) +{ + // 300 rows > block_size (256) to exercise a multi-block kernel launch + constexpr int N = 300; + std::vector vals(N); + std::iota(vals.begin(), vals.end(), 0); + cudf::test::fixed_width_column_wrapper col(vals.begin(), vals.end()); + cudf::table_view tbl{{col}}; + std::vector names{"v"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), N); + EXPECT_EQ(got->null_count(), 0); + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *got, "v", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} + +TEST_F(EncodeVariantTest, SlicedColumnInput) +{ + // Non-zero column offset: build a 5-row table and slice to rows [1, 4). + cudf::test::fixed_width_column_wrapper col_full{0, 10, 20, 30, 0}; + auto const col_sliced = cudf::slice(col_full, {1, 4})[0]; + cudf::table_view tbl{{col_sliced}}; + std::vector names{"x"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 3); + + cudf::test::fixed_width_column_wrapper expected_col{10, 20, 30}; + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *got, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, expected_col); +} From 0628356cab2884f372cf65ae4192f807e6deb8d9 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 5 Aug 2026 02:25:17 +0000 Subject: [PATCH 12/14] boolean support --- .../io/parquet/experimental/variant_encode.cu | 17 ++-- .../io/experimental/variant_encode_test.cpp | 78 +++++++++++++++++++ 2 files changed, 90 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index 13e1ade5def1..aa873b72cf57 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -42,6 +42,8 @@ namespace { // ────────────────────────────────────────────────────────────────────────────── constexpr uint8_t k_null_value = 0x00; +constexpr uint8_t k_bool_true_header = 0x04; // primitive type 1 (BOOLEAN_TRUE) +constexpr uint8_t k_bool_false_header = 0x08; // primitive type 2 (BOOLEAN_FALSE) constexpr uint8_t k_int8_header = 0x0c; constexpr uint8_t k_int16_header = 0x10; constexpr uint8_t k_int32_header = 0x14; @@ -71,6 +73,7 @@ __device__ int64_t field_encoded_size(column_device_view const& col, size_type r { if (col.type().id() == type_id::EMPTY || !col.is_valid(row)) { return 1; } switch (col.type().id()) { + case type_id::BOOL8: return 1; case type_id::INT8: return 2; case type_id::INT16: return 3; case type_id::INT32: return 5; @@ -93,6 +96,10 @@ __device__ size_type write_field_value(uint8_t* out, column_device_view const& c return 1; } switch (col.type().id()) { + case type_id::BOOL8: { + out[0] = col.element(row) ? k_bool_true_header : k_bool_false_header; + return 1; + } case type_id::INT8: { auto const v = col.element(row); out[0] = k_int8_header; @@ -477,11 +484,11 @@ std::unique_ptr encode_variant(cudf::table_view const& input, for (int i = 0; i < N; i++) { auto const id = input.column(i).type().id(); - CUDF_EXPECTS(id == type_id::EMPTY || id == type_id::INT8 || id == type_id::INT16 || - id == type_id::INT32 || id == type_id::INT64 || id == type_id::FLOAT32 || - id == type_id::FLOAT64 || id == type_id::STRING, - "encode_variant: unsupported column type — supported: EMPTY, INT8/16/32/64, " - "FLOAT32/64, STRING", + CUDF_EXPECTS(id == type_id::EMPTY || id == type_id::BOOL8 || id == type_id::INT8 || + id == type_id::INT16 || id == type_id::INT32 || id == type_id::INT64 || + id == type_id::FLOAT32 || id == type_id::FLOAT64 || id == type_id::STRING, + "encode_variant: unsupported column type — supported: EMPTY, BOOL8, " + "INT8/16/32/64, FLOAT32/64, STRING", std::invalid_argument); } diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index 6f78ed073423..b5ae0c5f42ce 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -30,6 +30,10 @@ namespace { inline std::vector enc_null() { return {0x00}; } +inline std::vector enc_bool_true() { return {0x04}; } + +inline std::vector enc_bool_false() { return {0x08}; } + inline std::vector enc_int8(int8_t v) { return {0x0c, static_cast(v)}; } inline std::vector enc_int16(int16_t v) @@ -632,3 +636,77 @@ TEST_F(EncodeVariantTest, SlicedColumnInput) *got, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, expected_col); } + +TEST_F(EncodeVariantTest, SingleBoolColumnBytes) +{ + // Verify boolean values produce the correct 1-byte header (no payload). + cudf::test::fixed_width_column_wrapper col{true, false}; + cudf::table_view tbl{{col}}; + std::vector names{"b"}; + + auto got = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(got->size(), 2); + EXPECT_EQ(got->null_count(), 0); + + auto const meta = build_metadata({"b"}); + auto const exp_val0 = build_object_value({enc_bool_true()}); + auto const exp_val1 = build_object_value({enc_bool_false()}); + + auto expected = make_variant_column({meta, meta}, {exp_val0, exp_val1}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(EncodeVariantTest, RoundtripBoolWithExtractVariant) +{ + cudf::test::fixed_width_column_wrapper col{true, false, true, false}; + cudf::table_view tbl{{col}}; + std::vector names{"flag"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *variant, "flag", cudf::data_type{cudf::type_id::BOOL8}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} + +TEST_F(EncodeVariantTest, RoundtripBoolNullsWithExtractVariant) +{ + // Null bool values should encode as VARIANT null and decode back as null BOOL8. + cudf::test::fixed_width_column_wrapper col({true, false, true}, {true, false, true}); + cudf::table_view tbl{{col}}; + std::vector names{"flag"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + EXPECT_EQ(variant->null_count(), 0); // struct rows never null in encode_variant + + auto decoded = cudf::io::parquet::experimental::extract_variant_field( + *variant, "flag", cudf::data_type{cudf::type_id::BOOL8}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, col); +} + +TEST_F(EncodeVariantTest, RoundtripBoolWithIntColumnMultiColumn) +{ + // Bool column alongside an int column — exercises multi-column sort and mixed types. + cudf::test::fixed_width_column_wrapper col_flag{true, false, true}; + cudf::test::fixed_width_column_wrapper col_id{1, 2, 3}; + cudf::table_view tbl{{col_flag, col_id}}; + std::vector names{"flag", "id"}; + + auto variant = + cudf::io::parquet::experimental::encode_variant(tbl, names, cudf::test::get_default_stream()); + + auto decoded_flag = cudf::io::parquet::experimental::extract_variant_field( + *variant, "flag", cudf::data_type{cudf::type_id::BOOL8}, cudf::test::get_default_stream()); + auto decoded_id = cudf::io::parquet::experimental::extract_variant_field( + *variant, "id", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded_flag, col_flag); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded_id, col_id); +} From 97a0374555d9372cc9cba0b4b10b67f16f30893e Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 5 Aug 2026 20:12:02 +0000 Subject: [PATCH 13/14] reviewers --- cpp/src/io/parquet/experimental/variant_encode.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index aa873b72cf57..89609cbbed2b 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -418,8 +418,8 @@ std::unique_ptr encode_strings_to_variant(cudf::strings_column_view cons CUDF_CUDA_TRY(cudaGetLastError()); } - auto [val_offsets_col, total_val_bytes] = cudf::strings::detail::make_offsets_child_column( - d_val_sizes.begin(), d_val_sizes.end(), stream, mr); + auto [val_offsets_col, total_val_bytes] = + cudf::detail::make_offsets_child_column(d_val_sizes.begin(), d_val_sizes.end(), stream, mr); // INT64 offsets are produced when total_bytes >= get_offset64_threshold() (== INT32_MAX). // Checking the column type is more robust than comparing total_val_bytes to INT32_MAX, // since the boundary case (total == INT32_MAX) produces INT64 but the numeric check passes. @@ -536,8 +536,8 @@ std::unique_ptr encode_variant(cudf::table_view const& input, CUDF_CUDA_TRY(cudaGetLastError()); } - auto [val_offsets_col, total_val_bytes] = cudf::strings::detail::make_offsets_child_column( - d_val_sizes.begin(), d_val_sizes.end(), stream, mr); + auto [val_offsets_col, total_val_bytes] = + cudf::detail::make_offsets_child_column(d_val_sizes.begin(), d_val_sizes.end(), stream, mr); CUDF_EXPECTS(val_offsets_col->type().id() == type_id::INT32, "VARIANT value bytes exceed 2 GiB limit", std::overflow_error); From 8d82b2fffe5ab789fc684e3325cd9fc5ec317144 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 6 Aug 2026 20:00:12 +0000 Subject: [PATCH 14/14] splitting pr --- cpp/include/cudf/io/experimental/variant.hpp | 22 --- .../io/parquet/experimental/variant_encode.cu | 161 ------------------ .../io/experimental/variant_encode_test.cpp | 157 ----------------- 3 files changed, 340 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 9bea68e21ccc..4bac662773f1 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -7,7 +7,6 @@ #include #include -#include #include #include #include @@ -113,27 +112,6 @@ namespace io::parquet::experimental { rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); -/** - * @brief Encode a STRING column as a scalar VARIANT column. - * - * Each non-null string is encoded as a VARIANT scalar value: a short_string blob (1 header byte + - * payload) for strings shorter than 64 bytes, or a long_string primitive blob (1 header + 4-byte - * LE length + payload) for longer strings. The metadata blob for every row is the minimal - * empty-dictionary encoding `{0x01, 0x00, 0x00}` (version 1, 0 keys). Null input rows produce - * null VARIANT struct rows. - * - * @param input STRING column to encode - * @param stream CUDA stream - * @param mr Device memory resource - * @return VARIANT column: `STRUCT, list>` (metadata child, value child) - * - * @throws cudf::logic_error if `input` is not a STRING column - */ -[[nodiscard]] std::unique_ptr encode_strings_to_variant( - cudf::strings_column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); - /** * @brief Encode a table as a VARIANT object column, one row per VARIANT. * diff --git a/cpp/src/io/parquet/experimental/variant_encode.cu b/cpp/src/io/parquet/experimental/variant_encode.cu index 89609cbbed2b..326d850177ed 100644 --- a/cpp/src/io/parquet/experimental/variant_encode.cu +++ b/cpp/src/io/parquet/experimental/variant_encode.cu @@ -13,7 +13,6 @@ #include #include #include -#include #include #include #include @@ -60,9 +59,6 @@ constexpr uint8_t k_short_string_basic_type = 0x01; // value_metadata = (value_header << 2) | basic_type::object(2) = 0x0e constexpr uint8_t k_object_value_metadata = 0x0e; -// Bytes in the constant empty-dictionary metadata blob {version=1, 0 keys} -constexpr size_type k_empty_meta_size = 3; - constexpr int block_size = 256; // ────────────────────────────────────────────────────────────────────────────── @@ -155,71 +151,6 @@ __device__ size_type write_field_value(uint8_t* out, column_device_view const& c } } -// ────────────────────────────────────────────────────────────────────────────── -// encode_strings_to_variant kernels -// ────────────────────────────────────────────────────────────────────────────── - -// Pass 1: compute per-row value blob sizes (metadata is always k_empty_meta_size = 3 bytes). -CUDF_KERNEL __launch_bounds__(block_size) void string_value_sizes_kernel( - column_device_view strings, device_span d_val_sizes) -{ - auto const num_rows = strings.size(); - auto const tid = cudf::detail::grid_1d::global_thread_id(); - auto const stride = cudf::detail::grid_1d::grid_stride(); - - for (auto row = tid; row < num_rows; row += stride) { - if (!strings.is_valid(row)) { - d_val_sizes[row] = 1; - continue; - } - auto const len = static_cast(strings.element(row).size_bytes()); - auto const enc_sz = len < 64 ? 1 + len : 5 + len; - d_val_sizes[row] = static_cast( - enc_sz <= static_cast(cuda::std::numeric_limits::max()) - ? enc_sz - : cuda::std::numeric_limits::max()); - } -} - -// Pass 2: write metadata and value blobs. -// Metadata is constant {0x01, 0x00, 0x00} at stride k_empty_meta_size bytes per row. -CUDF_KERNEL __launch_bounds__(block_size) void string_encode_write_kernel( - column_device_view strings, - device_span d_val_offsets, - uint8_t* d_val_buf, - uint8_t* d_meta_buf) -{ - auto const num_rows = strings.size(); - auto const tid = cudf::detail::grid_1d::global_thread_id(); - auto const stride = cudf::detail::grid_1d::grid_stride(); - - for (auto row = tid; row < num_rows; row += stride) { - // Constant metadata: {version=1, num_keys=0, sentinel_offset=0} - auto* mp = d_meta_buf + row * k_empty_meta_size; - mp[0] = 0x01; - mp[1] = 0x00; - mp[2] = 0x00; - - // Value - uint8_t* vp = d_val_buf + d_val_offsets[row]; - if (!strings.is_valid(row)) { - vp[0] = k_null_value; - continue; - } - auto const sv = strings.element(row); - auto const len = static_cast(sv.size_bytes()); - if (len < 64) { - vp[0] = static_cast(k_short_string_basic_type | (len << 2)); - cuda::std::memcpy(vp + 1, sv.data(), len); - } else { - vp[0] = k_long_string_header; - uint32_t const u32 = static_cast(len); - cuda::std::memcpy(vp + 1, &u32, 4); - cuda::std::memcpy(vp + 5, sv.data(), len); - } - } -} - // ────────────────────────────────────────────────────────────────────────────── // encode_variant kernels // ────────────────────────────────────────────────────────────────────────────── @@ -380,90 +311,6 @@ std::pair, std::unique_ptr> make_constant_list_b namespace detail { -// ────────────────────────────────────────────────────────────────────────────── -// encode_strings_to_variant -// ────────────────────────────────────────────────────────────────────────────── - -std::unique_ptr encode_strings_to_variant(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - auto const num_rows = input.size(); - - auto make_empty_list = [&] { - return make_lists_column( - 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); - }; - - if (num_rows == 0) { - std::vector> ch; - ch.push_back(make_empty_list()); - ch.push_back(make_empty_list()); - return make_structs_column(0, std::move(ch), 0, {}, stream, mr); - } - - auto input_dv = column_device_view::create(input.parent(), stream); - - // ── Metadata: constant k_empty_meta_size bytes per row ── - auto [meta_offsets_col, meta_data_col] = - make_constant_list_buffers(num_rows, k_empty_meta_size, stream, mr); - - // ── Value sizes ── - rmm::device_uvector d_val_sizes( - num_rows, stream, cudf::get_current_device_resource_ref()); - { - cudf::detail::grid_1d grid{num_rows, block_size}; - string_value_sizes_kernel<<>>(*input_dv, - d_val_sizes); - CUDF_CUDA_TRY(cudaGetLastError()); - } - - auto [val_offsets_col, total_val_bytes] = - cudf::detail::make_offsets_child_column(d_val_sizes.begin(), d_val_sizes.end(), stream, mr); - // INT64 offsets are produced when total_bytes >= get_offset64_threshold() (== INT32_MAX). - // Checking the column type is more robust than comparing total_val_bytes to INT32_MAX, - // since the boundary case (total == INT32_MAX) produces INT64 but the numeric check passes. - CUDF_EXPECTS(val_offsets_col->type().id() == type_id::INT32, - "VARIANT value bytes exceed 2 GiB limit", - std::overflow_error); - - auto val_data_col = make_numeric_column(data_type{type_id::UINT8}, - static_cast(total_val_bytes), - mask_state::UNALLOCATED, - stream, - mr); - - // ── Write pass (metadata + values) ── - { - device_span d_val_offsets{val_offsets_col->view().data(), - static_cast(num_rows + 1)}; - cudf::detail::grid_1d grid{num_rows, block_size}; - string_encode_write_kernel<<>>( - *input_dv, - d_val_offsets, - val_data_col->mutable_view().data(), - meta_data_col->mutable_view().data()); - CUDF_CUDA_TRY(cudaGetLastError()); - } - - // ── Null mask: propagate from input ── - size_type const null_count = input.null_count(); - rmm::device_buffer null_mask = - null_count > 0 ? cudf::detail::copy_bitmask(input.parent(), stream, mr) : rmm::device_buffer{}; - - // ── Assemble STRUCT, list> ── - auto meta_col = - make_lists_column(num_rows, std::move(meta_offsets_col), std::move(meta_data_col), 0, {}); - auto val_col = - make_lists_column(num_rows, std::move(val_offsets_col), std::move(val_data_col), 0, {}); - - std::vector> children; - children.push_back(std::move(meta_col)); - children.push_back(std::move(val_col)); - return make_structs_column( - num_rows, std::move(children), null_count, std::move(null_mask), stream, mr); -} - // ────────────────────────────────────────────────────────────────────────────── // encode_variant // ────────────────────────────────────────────────────────────────────────────── @@ -578,14 +425,6 @@ std::unique_ptr encode_variant(cudf::table_view const& input, } // namespace detail -std::unique_ptr encode_strings_to_variant(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - CUDF_FUNC_RANGE(); - return detail::encode_strings_to_variant(input, stream, mr); -} - std::unique_ptr encode_variant(cudf::table_view const& input, cudf::host_span column_names, rmm::cuda_stream_view stream, diff --git a/cpp/tests/io/experimental/variant_encode_test.cpp b/cpp/tests/io/experimental/variant_encode_test.cpp index b5ae0c5f42ce..96d173467a69 100644 --- a/cpp/tests/io/experimental/variant_encode_test.cpp +++ b/cpp/tests/io/experimental/variant_encode_test.cpp @@ -69,22 +69,6 @@ inline std::vector enc_short_string(std::string_view s) return out; } -inline std::vector enc_long_string(std::string_view s) -{ - auto const len = static_cast(s.size()); - std::vector out{0x40}; - for (int b = 0; b < 4; b++) { - out.push_back(static_cast((len >> (8 * b)) & 0xffu)); - } - out.insert(out.end(), s.begin(), s.end()); - return out; -} - -inline std::vector enc_string(std::string_view s) -{ - return s.size() < 64 ? enc_short_string(s) : enc_long_string(s); -} - // Build VARIANT metadata blob for a sorted list of key names (offset_size=1 assumed; keys must // have total byte length <= 255). inline std::vector build_metadata(std::vector const& sorted_keys) @@ -162,149 +146,8 @@ inline cudf::test::structs_column_wrapper make_variant_column( return cudf::test::structs_column_wrapper{std::move(children)}; } -// Return the metadata child column view (child 0) of a VARIANT struct column. -inline cudf::column_view meta_child(cudf::column_view const& variant_col) -{ - return cudf::structs_column_view{variant_col}.child(0); -} - -// Return the value child column view (child 1) of a VARIANT struct column. -inline cudf::column_view val_child(cudf::column_view const& variant_col) -{ - return cudf::structs_column_view{variant_col}.child(1); -} - } // namespace -// ────────────────────────────────────────────────────────────────────────────── -// encode_strings_to_variant tests -// ────────────────────────────────────────────────────────────────────────────── - -struct EncodeStringsToVariantTest : public cudf::test::BaseFixture {}; - -TEST_F(EncodeStringsToVariantTest, EmptyInput) -{ - cudf::test::strings_column_wrapper input{}; - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - EXPECT_EQ(got->type().id(), cudf::type_id::STRUCT); - EXPECT_EQ(got->size(), 0); - EXPECT_EQ(got->null_count(), 0); -} - -TEST_F(EncodeStringsToVariantTest, ShortString) -{ - cudf::test::strings_column_wrapper input{"hi"}; - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - EXPECT_EQ(got->size(), 1); - EXPECT_EQ(got->null_count(), 0); - - auto const expected_val = enc_short_string("hi"); - auto const expected_meta = std::vector{0x01, 0x00, 0x00}; - auto expected = make_variant_column({expected_meta}, {expected_val}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); -} - -TEST_F(EncodeStringsToVariantTest, LongString) -{ - // A string of exactly 64 bytes should use long_string encoding - std::string const s(64, 'x'); - cudf::test::strings_column_wrapper input{s}; - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - auto const expected_val = enc_long_string(s); - auto const expected_meta = std::vector{0x01, 0x00, 0x00}; - auto expected = make_variant_column({expected_meta}, {expected_val}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); -} - -TEST_F(EncodeStringsToVariantTest, NullInputProducesNullStructRow) -{ - cudf::test::strings_column_wrapper input({"hello", "", "world"}, {true, false, true}); - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - EXPECT_EQ(got->size(), 3); - EXPECT_EQ(got->null_count(), 1); -} - -TEST_F(EncodeStringsToVariantTest, AllNullInput) -{ - cudf::test::strings_column_wrapper input({"", "", ""}, {false, false, false}); - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - EXPECT_EQ(got->null_count(), 3); -} - -TEST_F(EncodeStringsToVariantTest, MultiRowMixedLengths) -{ - std::string const short_s = "abc"; // len 3 < 64 - std::string const long_s = std::string(100, 'z'); // len 100 >= 64 - - cudf::test::strings_column_wrapper input{short_s, long_s}; - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - EXPECT_EQ(got->size(), 2); - EXPECT_EQ(got->null_count(), 0); - - // Verify the value child has the right encoding for each row - auto const ev0 = enc_short_string(short_s); - auto const ev1 = enc_long_string(long_s); - auto expected = make_variant_column( - {std::vector{0x01, 0x00, 0x00}, std::vector{0x01, 0x00, 0x00}}, {ev0, ev1}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); -} - -TEST_F(EncodeStringsToVariantTest, RoundtripWithCastVariant) -{ - // encode then decode: cast_variant should recover the original strings - cudf::test::strings_column_wrapper input{"foo", "bar", "baz"}; - auto variant = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - auto decoded = cudf::io::parquet::experimental::cast_variant( - val_child(*variant), cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, input); -} - -TEST_F(EncodeStringsToVariantTest, RoundtripLongStringWithCastVariant) -{ - std::string const long_s = std::string(128, 'a'); - cudf::test::strings_column_wrapper input{long_s}; - auto variant = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{input}, cudf::test::get_default_stream()); - - auto decoded = cudf::io::parquet::experimental::cast_variant( - val_child(*variant), cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*decoded, input); -} - -TEST_F(EncodeStringsToVariantTest, SlicedInput) -{ - // Rows 0 and 3 are padding; the slice covers rows [1, 3) to exercise a non-zero column offset. - cudf::test::strings_column_wrapper full{"ignore", "hello", "world", "ignore"}; - auto const sliced = cudf::slice(full, {1, 3})[0]; - - auto got = cudf::io::parquet::experimental::encode_strings_to_variant( - cudf::strings_column_view{sliced}, cudf::test::get_default_stream()); - - EXPECT_EQ(got->size(), 2); - EXPECT_EQ(got->null_count(), 0); - - auto const meta = std::vector{0x01, 0x00, 0x00}; - auto const exp_v0 = enc_short_string("hello"); - auto const exp_v1 = enc_short_string("world"); - auto expected = make_variant_column({meta, meta}, {exp_v0, exp_v1}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); -} - // ────────────────────────────────────────────────────────────────────────────── // encode_variant tests // ──────────────────────────────────────────────────────────────────────────────