From a7f3578a1ecf34e8670985c81f47c6c60051ab94 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 31 Jul 2026 03:02:07 +0000 Subject: [PATCH 01/16] first draft of get_variant_id function --- cpp/include/cudf/io/experimental/variant.hpp | 22 ++ .../cudf/io/experimental/variant_spec.hpp | 24 ++ .../parquet/experimental/variant_extract.cu | 104 +++++ .../io/experimental/variant_extract_test.cpp | 354 ++++++++++++++++++ 4 files changed, 504 insertions(+) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 1b0719fad661..7059e24ffb8a 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 @@ -109,6 +110,27 @@ 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 Return the logical type of each VARIANT value blob in a `list` column. + * + * Physical integer widths INT8/INT16/INT32/INT64 all map to `long_value`; both string encodings + * (short and long) map to `string`. An encoded Variant null (NULLVAL) produces a valid + * `null_value` identifier — not a null output row. An input-null row produces an output-null row. + * An unrecognized or unknown header produces a null output row. + * + * @param values `list` column of VARIANT-encoded value bytes + * @param stream CUDA stream + * @param mr Device memory resource + * @return `INT32` column of `variant_logical_type` values cast to `int32_t`. A row is null when + * the input row is null or the value header carries an unrecognized type. + * + * @throws std::invalid_argument if `values` is not a `list` column + */ +[[nodiscard]] std::unique_ptr get_variant_type_id( + column_view const& values, + 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/include/cudf/io/experimental/variant_spec.hpp b/cpp/include/cudf/io/experimental/variant_spec.hpp index 6b71dc57385c..7cd0be4a730d 100644 --- a/cpp/include/cudf/io/experimental/variant_spec.hpp +++ b/cpp/include/cudf/io/experimental/variant_spec.hpp @@ -46,4 +46,28 @@ enum class variant_primitive_type : uint8_t { UUID = 20, }; +/** + * @brief Logical type of a VARIANT value as returned by get_variant_type_id. + * + * All four integer widths (INT8/INT16/INT32/INT64) map to long_value. Both string encodings + * (SHORT_STRING and LONG_STRING) map to string. The two timestamp-with-timezone encodings map to + * timestamp; the two timestamp-without-timezone encodings map to timestamp_ntz. + */ +enum class variant_logical_type : uint8_t { + object, + array, + null_value, + boolean, + long_value, + string, + double_value, + decimal, + date, + timestamp, + timestamp_ntz, + float_value, + binary, + uuid +}; + } // namespace cudf::io::parquet::experimental diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 5eefa1b47e79..22c9eed6db0c 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -752,6 +752,69 @@ struct cast_variant_fn { } }; +__device__ cuda::std::optional logical_type_of(device_span enc) +{ + if (enc.empty()) { return cuda::std::nullopt; } + auto const value_metadata = enc[0]; + auto const btype = decode_basic_type(value_metadata); + + if (btype == basic_type::SHORT_STRING) { return variant_logical_type::string; } + if (btype == basic_type::OBJECT) { return variant_logical_type::object; } + if (btype == basic_type::ARRAY) { return variant_logical_type::array; } + + switch (static_cast(variant_value_header(value_metadata))) { + case primitive_type::NULLVAL: return variant_logical_type::null_value; + case primitive_type::BOOLEAN_TRUE: + case primitive_type::BOOLEAN_FALSE: return variant_logical_type::boolean; + case primitive_type::INT8: + case primitive_type::INT16: + case primitive_type::INT32: + case primitive_type::INT64: return variant_logical_type::long_value; + case primitive_type::FLOAT64: return variant_logical_type::double_value; + case primitive_type::DECIMAL4: + case primitive_type::DECIMAL8: + case primitive_type::DECIMAL16: return variant_logical_type::decimal; + case primitive_type::DATE: return variant_logical_type::date; + case primitive_type::TIMESTAMP_MICROS: + case primitive_type::TIMESTAMP_NANOS: return variant_logical_type::timestamp; + case primitive_type::TIMESTAMP_NTZ_MICROS: + case primitive_type::TIMESTAMP_NTZ_NANOS: return variant_logical_type::timestamp_ntz; + case primitive_type::FLOAT32: return variant_logical_type::float_value; + case primitive_type::BINARY: return variant_logical_type::binary; + case primitive_type::LONG_STRING: return variant_logical_type::string; + case primitive_type::UUID: return variant_logical_type::uuid; + default: return cuda::std::nullopt; + } +} + +CUDF_KERNEL __launch_bounds__(block_size) void get_variant_type_id_kernel( + cudf::lists_column_device_view values, device_span d_output, bitmask_type* d_null_mask) +{ + auto const num_rows = static_cast(d_output.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 (!cudf::bit_is_set(d_null_mask, row)) { + d_output[row] = 0; + continue; + } + + auto const val_begin = values.offset_at(row); + auto const val_end = values.offset_at(row + 1); + device_span const val{values.child().data() + val_begin, + static_cast(val_end - val_begin)}; + + auto const ltype = logical_type_of(val); + if (ltype.has_value()) { + d_output[row] = static_cast(ltype.value()); + } else { + d_output[row] = 0; + cudf::clear_bit(d_null_mask, row); + } + } +} + std::unique_ptr build_path_column(cudf::host_span steps, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) @@ -908,6 +971,39 @@ std::unique_ptr cast_variant(column_view const& values, mr}); } +std::unique_ptr get_variant_type_id(column_view const& values, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + validate_variant_child(values); + size_type const num_rows = values.size(); + if (num_rows == 0) { return make_empty_column(data_type{type_id::INT32}); } + + auto val_device_view = column_device_view::create(values, stream); + cudf::lists_column_device_view val_lists_device_view(*val_device_view); + + auto null_mask = values.nullable() + ? cudf::detail::copy_bitmask(values, stream, mr) + : cudf::create_null_mask(num_rows, mask_state::ALL_VALID, stream, mr); + auto* d_null_mask = static_cast(null_mask.data()); + + rmm::device_buffer data{static_cast(num_rows) * sizeof(int32_t), stream, mr}; + + auto grid = cudf::detail::grid_1d{num_rows, block_size}; + get_variant_type_id_kernel<<>>( + val_lists_device_view, + {static_cast(data.data()), static_cast(num_rows)}, + d_null_mask); + CUDF_CUDA_TRY(cudaGetLastError()); + + auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); + return std::make_unique(data_type{type_id::INT32}, + num_rows, + std::move(data), + null_count > 0 ? std::move(null_mask) : rmm::device_buffer{}, + null_count); +} + } // namespace detail std::unique_ptr get_variant_field(column_view const& variant_column, @@ -928,6 +1024,14 @@ std::unique_ptr cast_variant(column_view const& values, return detail::cast_variant(values, desired_type, stream, mr); } +std::unique_ptr get_variant_type_id(column_view const& values, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::get_variant_type_id(values, stream, mr); +} + std::unique_ptr extract_variant_field(column_view const& variant_column, std::string_view path, data_type desired_type, diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 55f2291cbb76..0d3037d494f0 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include #include @@ -1301,3 +1302,356 @@ TEST_F(InvalidInputShapeTest, CastVariantRejectsMalformedInput) std::invalid_argument); } } + +// get_variant_type_id requires a list input; every other shape must be rejected with +// std::invalid_argument. +TEST_F(InvalidInputShapeTest, GetVariantTypeIdRejectsMalformedInput) +{ + auto stream = cudf::test::get_default_stream(); + + std::vector cases; + cases.push_back({"input is not a list", scalar_i32()}); + cases.push_back({"input list has wrong element type (not uint8)", list_i32({1, 2, 3})}); + + for (auto const& c : cases) { + SCOPED_TRACE(c.label); + EXPECT_THROW(static_cast( + cudf::io::parquet::experimental::get_variant_type_id(c.column->view(), stream)), + std::invalid_argument); + } +} + +// --------------------------------------------------------------------------- + +namespace { + +// Helper: run get_variant_type_id on the value child of an apache fixture. +template +std::unique_ptr apache_type_id(avf::fixture const& fixture) +{ + auto const stream = cudf::test::get_default_stream(); + auto col = make_apache_variant(fixture); + auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); + return cudf::io::parquet::experimental::get_variant_type_id(value, stream); +} + +// Build a list column from blobs with per-row validity. Rows where valid[i] is false are +// null at the list level (not an encoded Variant null — those are valid rows with a NULLVAL blob). +inline std::unique_ptr make_list_u8_nullable( + std::vector> const& blobs, std::vector const& valid) +{ + auto const n = static_cast(blobs.size()); + std::vector offs(n + 1, 0); + std::vector flat; + for (cudf::size_type i = 0; i < n; ++i) { + flat.insert(flat.end(), blobs[i].begin(), blobs[i].end()); + offs[i + 1] = static_cast(flat.size()); + } + auto off_col = + cudf::test::fixed_width_column_wrapper(offs.begin(), offs.end()).release(); + auto dat_col = + cudf::test::fixed_width_column_wrapper(flat.begin(), flat.end()).release(); + auto const null_count = + static_cast(std::count(valid.begin(), valid.end(), false)); + if (null_count == 0) { + return cudf::make_lists_column(n, std::move(off_col), std::move(dat_col), 0, {}); + } + auto const mask_bytes = cudf::bitmask_allocation_size_bytes(n); + std::vector host_mask(mask_bytes / sizeof(uint32_t), 0); + for (cudf::size_type i = 0; i < n; ++i) { + if (valid[i]) { host_mask[i / 32] |= uint32_t{1} << (i % 32); } + } + rmm::device_buffer d_mask(host_mask.data(), mask_bytes, cudf::test::get_default_stream()); + return cudf::make_lists_column( + n, std::move(off_col), std::move(dat_col), null_count, std::move(d_mask)); +} + +} // namespace + +struct GetVariantTypeIdTest : public cudf::test::BaseFixture {}; + +using LT = cudf::io::parquet::experimental::variant_logical_type; + +// --------------------------------------------------------------------------- +// Apache fixtures: one test per logical-type category. +// --------------------------------------------------------------------------- + +TEST_F(GetVariantTypeIdTest, NullValue) +{ + auto got = apache_type_id(avf::primitive_null); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::null_value)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, Boolean) +{ + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::boolean)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_boolean_true), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_boolean_false), expected); +} + +TEST_F(GetVariantTypeIdTest, LongValueAllIntWidths) +{ + // INT8, INT16, INT32, INT64 all map to long_value regardless of physical width. + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::long_value)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int8), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int16), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int32), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int64), expected); +} + +TEST_F(GetVariantTypeIdTest, StringBothEncodings) +{ + // SHORT_STRING and primitive LONG_STRING both map to string. + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::string)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::short_string), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_string), expected); +} + +TEST_F(GetVariantTypeIdTest, FloatTypes) +{ + { + auto got = apache_type_id(avf::primitive_float); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::float_value)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + { + auto got = apache_type_id(avf::primitive_double); + cudf::test::fixed_width_column_wrapper expected{ + static_cast(LT::double_value)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } +} + +TEST_F(GetVariantTypeIdTest, Decimal) +{ + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::decimal)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal4), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal8), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal16), expected); +} + +TEST_F(GetVariantTypeIdTest, Date) +{ + auto got = apache_type_id(avf::primitive_date); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::date)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, TimestampBothNanos) +{ + // TIMESTAMP_MICROS and TIMESTAMP_NANOS both map to timestamp. + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::timestamp)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestamp), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestamp_nanos), expected); +} + +TEST_F(GetVariantTypeIdTest, TimestampNtzBothNanos) +{ + // TIMESTAMP_NTZ_MICROS and TIMESTAMP_NTZ_NANOS both map to timestamp_ntz. + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::timestamp_ntz)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestampntz), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestampntz_nanos), expected); +} + +TEST_F(GetVariantTypeIdTest, Binary) +{ + auto got = apache_type_id(avf::primitive_binary); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::binary)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, Uuid) +{ + auto got = apache_type_id(avf::primitive_uuid); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::uuid)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, ObjectAndArray) +{ + { + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::object)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_primitive), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_nested), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_empty), expected); + } + { + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::array)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_primitive), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_nested), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_empty), expected); + } +} + +// --------------------------------------------------------------------------- +// Null and unknown-type behavior. +// --------------------------------------------------------------------------- + +TEST_F(GetVariantTypeIdTest, UnknownPhysicalTypeProducesNull) +{ + // TIME_NTZ_MICROS is a valid Variant physical type but has no logical-type mapping. + auto got = apache_type_id(avf::primitive_time); + ASSERT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(GetVariantTypeIdTest, InputNullRowPropagates) +{ + // A null row in the input list column propagates to the output. + auto const stream = cudf::test::get_default_stream(); + auto values = + make_list_u8_nullable({enc_int32(1), enc_int32(2), enc_int32(3)}, {true, false, true}); + + auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); + + cudf::test::fixed_width_column_wrapper expected( + {static_cast(LT::long_value), 0, static_cast(LT::long_value)}, + {true, false, true}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, EncodedNullIsNotInputNull) +{ + // An encoded Variant NULLVAL blob is a valid row whose type is null_value, not a null row. + auto const stream = cudf::test::get_default_stream(); + auto val = enc_null(); + cudf::test::lists_column_wrapper values(val.begin(), val.end()); + auto got = cudf::io::parquet::experimental::get_variant_type_id(values, stream); + + ASSERT_EQ(got->size(), 1); + EXPECT_EQ(got->null_count(), 0); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::null_value)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, EmptyValueBlobProducesNull) +{ + // An empty list row (zero bytes) has no header byte to decode → null. + auto const stream = cudf::test::get_default_stream(); + auto values = make_list_u8_nullable({enc_int32(1), {}, enc_int32(3)}, {true, true, true}); + + auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); + + cudf::test::fixed_width_column_wrapper expected( + {static_cast(LT::long_value), 0, static_cast(LT::long_value)}, + {true, false, true}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +// --------------------------------------------------------------------------- +// Multi-row and structural tests. +// --------------------------------------------------------------------------- + +TEST_F(GetVariantTypeIdTest, MixedTypesColumn) +{ + auto const stream = cudf::test::get_default_stream(); + + auto null_val = enc_null(); + auto bool_val = enc_bool(true); + auto int_val = enc_int64(999); + auto str_val = enc_short_string("hi"); + auto dbl_val = enc_float64(3.14); + + cudf::test::lists_column_wrapper values{ + {null_val.begin(), null_val.end()}, + {bool_val.begin(), bool_val.end()}, + {int_val.begin(), int_val.end()}, + {str_val.begin(), str_val.end()}, + {dbl_val.begin(), dbl_val.end()}, + }; + auto got = cudf::io::parquet::experimental::get_variant_type_id(values, stream); + + cudf::test::fixed_width_column_wrapper expected{ + static_cast(LT::null_value), + static_cast(LT::boolean), + static_cast(LT::long_value), + static_cast(LT::string), + static_cast(LT::double_value), + }; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, AllNullInputColumn) +{ + // All rows are null at the list level → all output rows are null. + auto const stream = cudf::test::get_default_stream(); + auto values = + make_list_u8_nullable({enc_int32(1), enc_int32(2), enc_int32(3)}, {false, false, false}); + + auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); + + ASSERT_EQ(got->size(), 3); + EXPECT_EQ(got->null_count(), 3); +} + +TEST_F(GetVariantTypeIdTest, EmptyInput) +{ + auto const stream = cudf::test::get_default_stream(); + auto const values = + cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); + + auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); + EXPECT_EQ(got->type().id(), cudf::type_id::INT32); + EXPECT_EQ(got->size(), 0); + EXPECT_EQ(got->null_count(), 0); +} + +TEST_F(GetVariantTypeIdTest, SlicedValuesColumn) +{ + // Verify that a sliced input produces correct results for the slice only. + auto const stream = cudf::test::get_default_stream(); + auto col = make_xyz_three_row_variant(); + auto const stream2 = cudf::test::get_default_stream(); + auto const value_child = cudf::structs_column_view{col}.get_sliced_child(1, stream2); + + // The xyz variant has object rows; slicing [1,3) gives 2 object rows. + auto const sliced_values = cudf::slice(value_child, {1, 3}).front(); + auto got = cudf::io::parquet::experimental::get_variant_type_id(sliced_values, stream); + + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::object), + static_cast(LT::object)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(GetVariantTypeIdTest, LargeMultiRowColumn) +{ + // 128 rows cycling through all types that get_variant_type_id can classify. + auto const stream = cudf::test::get_default_stream(); + + struct row_spec { + std::vector blob; + int32_t expected_id; + }; + + std::vector const types{ + {enc_null(), static_cast(LT::null_value)}, + {enc_bool(false), static_cast(LT::boolean)}, + {enc_int8(1), static_cast(LT::long_value)}, + {enc_int16(2), static_cast(LT::long_value)}, + {enc_int32(3), static_cast(LT::long_value)}, + {enc_int64(4), static_cast(LT::long_value)}, + {enc_float64(5.0), static_cast(LT::double_value)}, + {enc_short_string("x"), static_cast(LT::string)}, + {enc_long_string(std::string(70, 'z')), static_cast(LT::string)}, + }; + constexpr int num_rows = 128; + std::vector> blobs(num_rows); + std::vector expected_ids(num_rows); + for (int i = 0; i < num_rows; ++i) { + auto const& spec = types[i % types.size()]; + blobs[i] = spec.blob; + expected_ids[i] = spec.expected_id; + } + + auto values = make_list_u8_nullable(blobs, std::vector(num_rows, true)); + auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); + + cudf::test::fixed_width_column_wrapper expected(expected_ids.begin(), + expected_ids.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} From d169aa4da6a486d9359ec465406db04c58a9e416 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 5 Aug 2026 01:50:50 +0000 Subject: [PATCH 02/16] reviews --- cpp/tests/io/experimental/variant_extract_test.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 0d3037d494f0..d359bc0113e2 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1620,7 +1620,8 @@ TEST_F(GetVariantTypeIdTest, SlicedValuesColumn) TEST_F(GetVariantTypeIdTest, LargeMultiRowColumn) { - // 128 rows cycling through all types that get_variant_type_id can classify. + // 600 rows cycling through all types that get_variant_type_id can classify. + // 600 > 512 (typical block size) so the kernel exercises the multi-block grid-stride path. auto const stream = cudf::test::get_default_stream(); struct row_spec { @@ -1639,7 +1640,7 @@ TEST_F(GetVariantTypeIdTest, LargeMultiRowColumn) {enc_short_string("x"), static_cast(LT::string)}, {enc_long_string(std::string(70, 'z')), static_cast(LT::string)}, }; - constexpr int num_rows = 128; + constexpr int num_rows = 600; std::vector> blobs(num_rows); std::vector expected_ids(num_rows); for (int i = 0; i < num_rows; ++i) { From 3f1e0a84d6af34d786335b18b6cd6a5af2bd6ed9 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 5 Aug 2026 14:00:36 -0500 Subject: [PATCH 03/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 6b2880ffd4a2..bd6a0b810e1d 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1401,7 +1401,6 @@ TEST_F(InvalidInputShapeTest, GetVariantTypeIdRejectsMalformedInput) } } -// --------------------------------------------------------------------------- namespace { From e3930a490a6a54f8b0ec39174b6673a4fc0c34bc Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 5 Aug 2026 14:00:42 -0500 Subject: [PATCH 04/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index bd6a0b810e1d..b9773cf4683e 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1449,7 +1449,7 @@ inline std::unique_ptr make_list_u8_nullable( struct GetVariantTypeIdTest : public cudf::test::BaseFixture {}; -using LT = cudf::io::parquet::experimental::variant_logical_type; +using cudf::io::parquet::experimental::variant_logical_type; // --------------------------------------------------------------------------- // Apache fixtures: one test per logical-type category. From 512ffdc95025888e7437169d6a48eece815fc621 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 5 Aug 2026 19:50:04 +0000 Subject: [PATCH 05/16] addressing reviews --- .../cudf/io/experimental/variant_spec.hpp | 4 +- .../parquet/experimental/variant_extract.cu | 47 +++++-------------- .../io/experimental/variant_extract_test.cpp | 15 ++++-- 3 files changed, 28 insertions(+), 38 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant_spec.hpp b/cpp/include/cudf/io/experimental/variant_spec.hpp index 7cd0be4a730d..028add8eaeac 100644 --- a/cpp/include/cudf/io/experimental/variant_spec.hpp +++ b/cpp/include/cudf/io/experimental/variant_spec.hpp @@ -52,6 +52,7 @@ enum class variant_primitive_type : uint8_t { * All four integer widths (INT8/INT16/INT32/INT64) map to long_value. Both string encodings * (SHORT_STRING and LONG_STRING) map to string. The two timestamp-with-timezone encodings map to * timestamp; the two timestamp-without-timezone encodings map to timestamp_ntz. + * TIME_NTZ_MICROS maps to time_ntz. */ enum class variant_logical_type : uint8_t { object, @@ -67,7 +68,8 @@ enum class variant_logical_type : uint8_t { timestamp_ntz, float_value, binary, - uuid + uuid, + time_ntz }; } // namespace cudf::io::parquet::experimental diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index beaa5df7fe0f..13277b9f33e0 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -822,39 +822,12 @@ __device__ cuda::std::optional logical_type_of(device_span case primitive_type::FLOAT32: return variant_logical_type::float_value; case primitive_type::BINARY: return variant_logical_type::binary; case primitive_type::LONG_STRING: return variant_logical_type::string; + case primitive_type::TIME_NTZ_MICROS: return variant_logical_type::time_ntz; case primitive_type::UUID: return variant_logical_type::uuid; default: return cuda::std::nullopt; } } -CUDF_KERNEL __launch_bounds__(block_size) void get_variant_type_id_kernel( - cudf::lists_column_device_view values, device_span d_output, bitmask_type* d_null_mask) -{ - auto const num_rows = static_cast(d_output.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 (!cudf::bit_is_set(d_null_mask, row)) { - d_output[row] = 0; - continue; - } - - auto const val_begin = values.offset_at(row); - auto const val_end = values.offset_at(row + 1); - device_span const val{values.child().data() + val_begin, - static_cast(val_end - val_begin)}; - - auto const ltype = logical_type_of(val); - if (ltype.has_value()) { - d_output[row] = static_cast(ltype.value()); - } else { - d_output[row] = 0; - cudf::clear_bit(d_null_mask, row); - } - } -} - std::unique_ptr build_path_column(cudf::host_span steps, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) @@ -1029,12 +1002,18 @@ std::unique_ptr get_variant_type_id(column_view const& values, rmm::device_buffer data{static_cast(num_rows) * sizeof(int32_t), stream, mr}; - auto grid = cudf::detail::grid_1d{num_rows, block_size}; - get_variant_type_id_kernel<<>>( - val_lists_device_view, - {static_cast(data.data()), static_cast(num_rows)}, - d_null_mask); - CUDF_CUDA_TRY(cudaGetLastError()); + thrust::transform( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator(0), + cuda::counting_iterator(num_rows), + static_cast(data.data()), + [values = val_lists_device_view, d_null_mask] __device__(size_type row) -> int32_t { + if (!cudf::bit_is_set(d_null_mask, row)) { return 0; } + auto const ltype = logical_type_of(list_row_span(values, row)); + if (ltype.has_value()) { return static_cast(ltype.value()); } + cudf::clear_bit(d_null_mask, row); + return 0; + }); auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); return std::make_unique(data_type{type_id::INT32}, diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index b9773cf4683e..0d9b0c520ad6 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1401,7 +1401,6 @@ TEST_F(InvalidInputShapeTest, GetVariantTypeIdRejectsMalformedInput) } } - namespace { // Helper: run get_variant_type_id on the value child of an apache fixture. @@ -1450,6 +1449,7 @@ inline std::unique_ptr make_list_u8_nullable( struct GetVariantTypeIdTest : public cudf::test::BaseFixture {}; using cudf::io::parquet::experimental::variant_logical_type; +using LT = variant_logical_type; // --------------------------------------------------------------------------- // Apache fixtures: one test per logical-type category. @@ -1550,6 +1550,13 @@ TEST_F(GetVariantTypeIdTest, Uuid) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } +TEST_F(GetVariantTypeIdTest, TimeNtz) +{ + auto got = apache_type_id(avf::primitive_time); + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::time_ntz)}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + TEST_F(GetVariantTypeIdTest, ObjectAndArray) { { @@ -1573,8 +1580,10 @@ TEST_F(GetVariantTypeIdTest, ObjectAndArray) TEST_F(GetVariantTypeIdTest, UnknownPhysicalTypeProducesNull) { - // TIME_NTZ_MICROS is a valid Variant physical type but has no logical-type mapping. - auto got = apache_type_id(avf::primitive_time); + // Primitive header byte 0xFC = (63 << 2) | 0: type_id 63 is not in the spec. + auto const stream = cudf::test::get_default_stream(); + auto values = make_list_u8_nullable({{0xFC}}, {true}); + auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); } From e5145cff1bff060e0c3b6bfc84e9bb88ba28f3c6 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:38:54 -0500 Subject: [PATCH 06/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 0d9b0c520ad6..8318ce52f4cc 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1405,7 +1405,7 @@ namespace { // Helper: run get_variant_type_id on the value child of an apache fixture. template -std::unique_ptr apache_type_id(avf::fixture const& fixture) +[[nodiscard]] std::unique_ptr apache_type_id(avf::fixture const& fixture) { auto const stream = cudf::test::get_default_stream(); auto col = make_apache_variant(fixture); From 2b18cf54dc10e4f91a9a8eeb193cdabf113608a0 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:39:04 -0500 Subject: [PATCH 07/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 8318ce52f4cc..5728d4d5c8af 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1415,7 +1415,7 @@ template // Build a list column from blobs with per-row validity. Rows where valid[i] is false are // null at the list level (not an encoded Variant null — those are valid rows with a NULLVAL blob). -inline std::unique_ptr make_list_u8_nullable( +[[nodiscard]] std::unique_ptr make_list_u8_nullable( std::vector> const& blobs, std::vector const& valid) { auto const n = static_cast(blobs.size()); From d0db423f1fd08afce469399a578aea5ce4d84dd5 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:39:17 -0500 Subject: [PATCH 08/16] Update cpp/include/cudf/io/experimental/variant_spec.hpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/include/cudf/io/experimental/variant_spec.hpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant_spec.hpp b/cpp/include/cudf/io/experimental/variant_spec.hpp index 028add8eaeac..e4f71605bd78 100644 --- a/cpp/include/cudf/io/experimental/variant_spec.hpp +++ b/cpp/include/cudf/io/experimental/variant_spec.hpp @@ -49,10 +49,6 @@ enum class variant_primitive_type : uint8_t { /** * @brief Logical type of a VARIANT value as returned by get_variant_type_id. * - * All four integer widths (INT8/INT16/INT32/INT64) map to long_value. Both string encodings - * (SHORT_STRING and LONG_STRING) map to string. The two timestamp-with-timezone encodings map to - * timestamp; the two timestamp-without-timezone encodings map to timestamp_ntz. - * TIME_NTZ_MICROS maps to time_ntz. */ enum class variant_logical_type : uint8_t { object, From 1445a9dbc0cedcab1ca6db379d665640be9b162c Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:39:48 -0500 Subject: [PATCH 09/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 5728d4d5c8af..10bf002e6010 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1448,8 +1448,7 @@ template struct GetVariantTypeIdTest : public cudf::test::BaseFixture {}; -using cudf::io::parquet::experimental::variant_logical_type; -using LT = variant_logical_type; +using vlt = cudf::io::parquet::experimental::variant_logical_type; // --------------------------------------------------------------------------- // Apache fixtures: one test per logical-type category. From b61ea8aaffcbf20aea3cae98782c4277a8c1c8a7 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:40:06 -0500 Subject: [PATCH 10/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 10bf002e6010..1509ae2fabff 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1457,7 +1457,7 @@ using vlt = cudf::io::parquet::experimental::variant_logical_type; TEST_F(GetVariantTypeIdTest, NullValue) { auto got = apache_type_id(avf::primitive_null); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::null_value)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(vlt::null_value)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } From ff660ee36ebfd21c9113999ed5271589136d7202 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:40:36 -0500 Subject: [PATCH 11/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 1509ae2fabff..902d326f0c40 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1693,7 +1693,6 @@ TEST_F(GetVariantTypeIdTest, SlicedValuesColumn) // Verify that a sliced input produces correct results for the slice only. auto const stream = cudf::test::get_default_stream(); auto col = make_xyz_three_row_variant(); - auto const stream2 = cudf::test::get_default_stream(); auto const value_child = cudf::structs_column_view{col}.get_sliced_child(1, stream2); // The xyz variant has object rows; slicing [1,3) gives 2 object rows. From c91b9057147aa7d47bddaf112e5e1096406eb30f Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:46:30 -0500 Subject: [PATCH 12/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 902d326f0c40..886b5188c29f 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1693,7 +1693,7 @@ TEST_F(GetVariantTypeIdTest, SlicedValuesColumn) // Verify that a sliced input produces correct results for the slice only. auto const stream = cudf::test::get_default_stream(); auto col = make_xyz_three_row_variant(); - auto const value_child = cudf::structs_column_view{col}.get_sliced_child(1, stream2); + auto const value_child = cudf::structs_column_view{col}.get_sliced_child(1, stream); // The xyz variant has object rows; slicing [1,3) gives 2 object rows. auto const sliced_values = cudf::slice(value_child, {1, 3}).front(); From 7b646248bdcd9662438a263cad25ca00181960cb Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 16:46:43 -0500 Subject: [PATCH 13/16] Update cpp/tests/io/experimental/variant_extract_test.cpp Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com> --- cpp/tests/io/experimental/variant_extract_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 886b5188c29f..324ab474a3c2 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1419,7 +1419,7 @@ template std::vector> const& blobs, std::vector const& valid) { auto const n = static_cast(blobs.size()); - std::vector offs(n + 1, 0); + std::vector offsets(n + 1, 0); std::vector flat; for (cudf::size_type i = 0; i < n; ++i) { flat.insert(flat.end(), blobs[i].begin(), blobs[i].end()); From 5d9c0bbadd7c3f8581253cc5b0fc08bf77e7a4ec Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 7 Aug 2026 22:08:26 +0000 Subject: [PATCH 14/16] reviews --- cpp/include/cudf/io/experimental/variant.hpp | 11 +- .../cudf/io/experimental/variant_spec.hpp | 30 ++--- .../parquet/experimental/variant_extract.cu | 35 +++--- .../io/experimental/variant_extract_test.cpp | 111 ++++++++---------- 4 files changed, 91 insertions(+), 96 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 8738e8a4b86f..53729920efa7 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -113,16 +113,15 @@ namespace io::parquet::experimental { /** * @brief Return the logical type of each VARIANT value blob in a `list` column. * - * Physical integer widths INT8/INT16/INT32/INT64 all map to `long_value`; both string encodings - * (short and long) map to `string`. An encoded Variant null (NULLVAL) produces a valid - * `null_value` identifier — not a null output row. An input-null row produces an output-null row. - * An unrecognized or unknown header produces a null output row. + * Classifies only the value_metadata header byte; does not validate the remaining payload. + * A recognized header returns its logical type even when the payload is truncated. A null output + * row is produced when the input row is null, the blob is empty, or the header carries an + * unrecognized type. An encoded Variant null (NULLVAL) produces a valid `NULL_VALUE` row. * * @param values `list` column of VARIANT-encoded value bytes * @param stream CUDA stream * @param mr Device memory resource - * @return `INT32` column of `variant_logical_type` values cast to `int32_t`. A row is null when - * the input row is null or the value header carries an unrecognized type. + * @return `INT32` column of `variant_logical_type` values cast to `int32_t` * * @throws std::invalid_argument if `values` is not a `list` column */ diff --git a/cpp/include/cudf/io/experimental/variant_spec.hpp b/cpp/include/cudf/io/experimental/variant_spec.hpp index e4f71605bd78..b34da3564556 100644 --- a/cpp/include/cudf/io/experimental/variant_spec.hpp +++ b/cpp/include/cudf/io/experimental/variant_spec.hpp @@ -51,21 +51,21 @@ enum class variant_primitive_type : uint8_t { * */ enum class variant_logical_type : uint8_t { - object, - array, - null_value, - boolean, - long_value, - string, - double_value, - decimal, - date, - timestamp, - timestamp_ntz, - float_value, - binary, - uuid, - time_ntz + OBJECT = 0, + ARRAY = 1, + NULL_VALUE = 2, + BOOLEAN = 3, + LONG_VALUE = 4, + STRING = 5, + DOUBLE_VALUE = 6, + DECIMAL = 7, + DATE = 8, + TIMESTAMP = 9, + TIMESTAMP_NTZ = 10, + FLOAT_VALUE = 11, + BINARY = 12, + UUID = 13, + TIME_NTZ = 14, }; } // namespace cudf::io::parquet::experimental diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 13277b9f33e0..24e5a7602ba6 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -792,38 +792,41 @@ struct cast_variant_fn { } }; +// Classifies only the first (value_metadata) byte of enc; does not validate the remaining payload. +// A recognized header returns its logical type even when the payload is truncated. +// Returns nullopt for an empty blob or an unrecognized primitive type ID. __device__ cuda::std::optional logical_type_of(device_span enc) { if (enc.empty()) { return cuda::std::nullopt; } auto const value_metadata = enc[0]; auto const btype = decode_basic_type(value_metadata); - if (btype == basic_type::SHORT_STRING) { return variant_logical_type::string; } - if (btype == basic_type::OBJECT) { return variant_logical_type::object; } - if (btype == basic_type::ARRAY) { return variant_logical_type::array; } + if (btype == basic_type::SHORT_STRING) { return variant_logical_type::STRING; } + if (btype == basic_type::OBJECT) { return variant_logical_type::OBJECT; } + if (btype == basic_type::ARRAY) { return variant_logical_type::ARRAY; } switch (static_cast(variant_value_header(value_metadata))) { - case primitive_type::NULLVAL: return variant_logical_type::null_value; + case primitive_type::NULLVAL: return variant_logical_type::NULL_VALUE; case primitive_type::BOOLEAN_TRUE: - case primitive_type::BOOLEAN_FALSE: return variant_logical_type::boolean; + case primitive_type::BOOLEAN_FALSE: return variant_logical_type::BOOLEAN; case primitive_type::INT8: case primitive_type::INT16: case primitive_type::INT32: - case primitive_type::INT64: return variant_logical_type::long_value; - case primitive_type::FLOAT64: return variant_logical_type::double_value; + case primitive_type::INT64: return variant_logical_type::LONG_VALUE; + case primitive_type::FLOAT64: return variant_logical_type::DOUBLE_VALUE; case primitive_type::DECIMAL4: case primitive_type::DECIMAL8: - case primitive_type::DECIMAL16: return variant_logical_type::decimal; - case primitive_type::DATE: return variant_logical_type::date; + case primitive_type::DECIMAL16: return variant_logical_type::DECIMAL; + case primitive_type::DATE: return variant_logical_type::DATE; case primitive_type::TIMESTAMP_MICROS: - case primitive_type::TIMESTAMP_NANOS: return variant_logical_type::timestamp; + case primitive_type::TIMESTAMP_NANOS: return variant_logical_type::TIMESTAMP; case primitive_type::TIMESTAMP_NTZ_MICROS: - case primitive_type::TIMESTAMP_NTZ_NANOS: return variant_logical_type::timestamp_ntz; - case primitive_type::FLOAT32: return variant_logical_type::float_value; - case primitive_type::BINARY: return variant_logical_type::binary; - case primitive_type::LONG_STRING: return variant_logical_type::string; - case primitive_type::TIME_NTZ_MICROS: return variant_logical_type::time_ntz; - case primitive_type::UUID: return variant_logical_type::uuid; + case primitive_type::TIMESTAMP_NTZ_NANOS: return variant_logical_type::TIMESTAMP_NTZ; + case primitive_type::FLOAT32: return variant_logical_type::FLOAT_VALUE; + case primitive_type::BINARY: return variant_logical_type::BINARY; + case primitive_type::LONG_STRING: return variant_logical_type::STRING; + case primitive_type::TIME_NTZ_MICROS: return variant_logical_type::TIME_NTZ; + case primitive_type::UUID: return variant_logical_type::UUID; default: return cuda::std::nullopt; } } diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 324ab474a3c2..0b1b8ff82c73 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1416,39 +1416,29 @@ template // Build a list column from blobs with per-row validity. Rows where valid[i] is false are // null at the list level (not an encoded Variant null — those are valid rows with a NULLVAL blob). [[nodiscard]] std::unique_ptr make_list_u8_nullable( - std::vector> const& blobs, std::vector const& valid) + cudf::host_span const> blobs, std::vector const& valid) { - auto const n = static_cast(blobs.size()); - std::vector offsets(n + 1, 0); + auto const num_rows = static_cast(blobs.size()); + std::vector offsets(num_rows + 1, 0); std::vector flat; - for (cudf::size_type i = 0; i < n; ++i) { + for (cudf::size_type i = 0; i < num_rows; ++i) { flat.insert(flat.end(), blobs[i].begin(), blobs[i].end()); - offs[i + 1] = static_cast(flat.size()); + offsets[i + 1] = static_cast(flat.size()); } auto off_col = - cudf::test::fixed_width_column_wrapper(offs.begin(), offs.end()).release(); + cudf::test::fixed_width_column_wrapper(offsets.begin(), offsets.end()).release(); auto dat_col = cudf::test::fixed_width_column_wrapper(flat.begin(), flat.end()).release(); - auto const null_count = - static_cast(std::count(valid.begin(), valid.end(), false)); - if (null_count == 0) { - return cudf::make_lists_column(n, std::move(off_col), std::move(dat_col), 0, {}); - } - auto const mask_bytes = cudf::bitmask_allocation_size_bytes(n); - std::vector host_mask(mask_bytes / sizeof(uint32_t), 0); - for (cudf::size_type i = 0; i < n; ++i) { - if (valid[i]) { host_mask[i / 32] |= uint32_t{1} << (i % 32); } - } - rmm::device_buffer d_mask(host_mask.data(), mask_bytes, cudf::test::get_default_stream()); + auto [d_mask, null_count] = cudf::test::detail::make_null_mask(valid.begin(), valid.end()); return cudf::make_lists_column( - n, std::move(off_col), std::move(dat_col), null_count, std::move(d_mask)); + num_rows, std::move(off_col), std::move(dat_col), null_count, std::move(d_mask)); } } // namespace struct GetVariantTypeIdTest : public cudf::test::BaseFixture {}; -using vlt = cudf::io::parquet::experimental::variant_logical_type; +using LT = cudf::io::parquet::experimental::variant_logical_type; // --------------------------------------------------------------------------- // Apache fixtures: one test per logical-type category. @@ -1457,13 +1447,13 @@ using vlt = cudf::io::parquet::experimental::variant_logical_type; TEST_F(GetVariantTypeIdTest, NullValue) { auto got = apache_type_id(avf::primitive_null); - cudf::test::fixed_width_column_wrapper expected{static_cast(vlt::null_value)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::NULL_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, Boolean) { - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::boolean)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::BOOLEAN)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_boolean_true), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_boolean_false), expected); } @@ -1472,7 +1462,7 @@ TEST_F(GetVariantTypeIdTest, LongValueAllIntWidths) { // INT8, INT16, INT32, INT64 all map to long_value regardless of physical width. cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::long_value)}; + static_cast(LT::LONG_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int8), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int16), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int32), expected); @@ -1482,7 +1472,7 @@ TEST_F(GetVariantTypeIdTest, LongValueAllIntWidths) TEST_F(GetVariantTypeIdTest, StringBothEncodings) { // SHORT_STRING and primitive LONG_STRING both map to string. - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::string)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::STRING)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::short_string), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_string), expected); } @@ -1491,20 +1481,20 @@ TEST_F(GetVariantTypeIdTest, FloatTypes) { { auto got = apache_type_id(avf::primitive_float); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::float_value)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::FLOAT_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = apache_type_id(avf::primitive_double); cudf::test::fixed_width_column_wrapper expected{ - static_cast(LT::double_value)}; + static_cast(LT::DOUBLE_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } TEST_F(GetVariantTypeIdTest, Decimal) { - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::decimal)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::DECIMAL)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal4), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal8), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal16), expected); @@ -1513,7 +1503,7 @@ TEST_F(GetVariantTypeIdTest, Decimal) TEST_F(GetVariantTypeIdTest, Date) { auto got = apache_type_id(avf::primitive_date); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::date)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::DATE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1521,7 +1511,7 @@ TEST_F(GetVariantTypeIdTest, TimestampBothNanos) { // TIMESTAMP_MICROS and TIMESTAMP_NANOS both map to timestamp. cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::timestamp)}; + static_cast(LT::TIMESTAMP)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestamp), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestamp_nanos), expected); } @@ -1530,7 +1520,7 @@ TEST_F(GetVariantTypeIdTest, TimestampNtzBothNanos) { // TIMESTAMP_NTZ_MICROS and TIMESTAMP_NTZ_NANOS both map to timestamp_ntz. cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::timestamp_ntz)}; + static_cast(LT::TIMESTAMP_NTZ)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestampntz), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestampntz_nanos), expected); } @@ -1538,21 +1528,21 @@ TEST_F(GetVariantTypeIdTest, TimestampNtzBothNanos) TEST_F(GetVariantTypeIdTest, Binary) { auto got = apache_type_id(avf::primitive_binary); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::binary)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::BINARY)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, Uuid) { auto got = apache_type_id(avf::primitive_uuid); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::uuid)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::UUID)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, TimeNtz) { auto got = apache_type_id(avf::primitive_time); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::time_ntz)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::TIME_NTZ)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1560,13 +1550,13 @@ TEST_F(GetVariantTypeIdTest, ObjectAndArray) { { cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::object)}; + static_cast(LT::OBJECT)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_primitive), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_nested), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_empty), expected); } { - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::array)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::ARRAY)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_primitive), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_nested), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_empty), expected); @@ -1581,7 +1571,7 @@ TEST_F(GetVariantTypeIdTest, UnknownPhysicalTypeProducesNull) { // Primitive header byte 0xFC = (63 << 2) | 0: type_id 63 is not in the spec. auto const stream = cudf::test::get_default_stream(); - auto values = make_list_u8_nullable({{0xFC}}, {true}); + auto values = make_list_u8_nullable(std::vector>{{0xFC}}, {true}); auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); @@ -1591,13 +1581,14 @@ TEST_F(GetVariantTypeIdTest, InputNullRowPropagates) { // A null row in the input list column propagates to the output. auto const stream = cudf::test::get_default_stream(); - auto values = - make_list_u8_nullable({enc_int32(1), enc_int32(2), enc_int32(3)}, {true, false, true}); + auto values = make_list_u8_nullable( + std::vector>{enc_int32(1), enc_int32(2), enc_int32(3)}, + {true, false, true}); auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); cudf::test::fixed_width_column_wrapper expected( - {static_cast(LT::long_value), 0, static_cast(LT::long_value)}, + {static_cast(LT::LONG_VALUE), 0, static_cast(LT::LONG_VALUE)}, {true, false, true}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1612,7 +1603,7 @@ TEST_F(GetVariantTypeIdTest, EncodedNullIsNotInputNull) ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 0); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::null_value)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::NULL_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1620,12 +1611,13 @@ TEST_F(GetVariantTypeIdTest, EmptyValueBlobProducesNull) { // An empty list row (zero bytes) has no header byte to decode → null. auto const stream = cudf::test::get_default_stream(); - auto values = make_list_u8_nullable({enc_int32(1), {}, enc_int32(3)}, {true, true, true}); + auto values = make_list_u8_nullable( + std::vector>{enc_int32(1), {}, enc_int32(3)}, {true, true, true}); auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); cudf::test::fixed_width_column_wrapper expected( - {static_cast(LT::long_value), 0, static_cast(LT::long_value)}, + {static_cast(LT::LONG_VALUE), 0, static_cast(LT::LONG_VALUE)}, {true, false, true}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1654,11 +1646,11 @@ TEST_F(GetVariantTypeIdTest, MixedTypesColumn) auto got = cudf::io::parquet::experimental::get_variant_type_id(values, stream); cudf::test::fixed_width_column_wrapper expected{ - static_cast(LT::null_value), - static_cast(LT::boolean), - static_cast(LT::long_value), - static_cast(LT::string), - static_cast(LT::double_value), + static_cast(LT::NULL_VALUE), + static_cast(LT::BOOLEAN), + static_cast(LT::LONG_VALUE), + static_cast(LT::STRING), + static_cast(LT::DOUBLE_VALUE), }; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1667,8 +1659,9 @@ TEST_F(GetVariantTypeIdTest, AllNullInputColumn) { // All rows are null at the list level → all output rows are null. auto const stream = cudf::test::get_default_stream(); - auto values = - make_list_u8_nullable({enc_int32(1), enc_int32(2), enc_int32(3)}, {false, false, false}); + auto values = make_list_u8_nullable( + std::vector>{enc_int32(1), enc_int32(2), enc_int32(3)}, + {false, false, false}); auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); @@ -1699,8 +1692,8 @@ TEST_F(GetVariantTypeIdTest, SlicedValuesColumn) auto const sliced_values = cudf::slice(value_child, {1, 3}).front(); auto got = cudf::io::parquet::experimental::get_variant_type_id(sliced_values, stream); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::object), - static_cast(LT::object)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::OBJECT), + static_cast(LT::OBJECT)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1716,15 +1709,15 @@ TEST_F(GetVariantTypeIdTest, LargeMultiRowColumn) }; std::vector const types{ - {enc_null(), static_cast(LT::null_value)}, - {enc_bool(false), static_cast(LT::boolean)}, - {enc_int8(1), static_cast(LT::long_value)}, - {enc_int16(2), static_cast(LT::long_value)}, - {enc_int32(3), static_cast(LT::long_value)}, - {enc_int64(4), static_cast(LT::long_value)}, - {enc_float64(5.0), static_cast(LT::double_value)}, - {enc_short_string("x"), static_cast(LT::string)}, - {enc_long_string(std::string(70, 'z')), static_cast(LT::string)}, + {enc_null(), static_cast(LT::NULL_VALUE)}, + {enc_bool(false), static_cast(LT::BOOLEAN)}, + {enc_int8(1), static_cast(LT::LONG_VALUE)}, + {enc_int16(2), static_cast(LT::LONG_VALUE)}, + {enc_int32(3), static_cast(LT::LONG_VALUE)}, + {enc_int64(4), static_cast(LT::LONG_VALUE)}, + {enc_float64(5.0), static_cast(LT::DOUBLE_VALUE)}, + {enc_short_string("x"), static_cast(LT::STRING)}, + {enc_long_string(std::string(70, 'z')), static_cast(LT::STRING)}, }; constexpr int num_rows = 600; std::vector> blobs(num_rows); From 45033053e4f6f674c89ff9f6440c271dd3d702b8 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Sat, 8 Aug 2026 02:04:57 +0000 Subject: [PATCH 15/16] nits --- cpp/include/cudf/io/experimental/variant_spec.hpp | 1 - .../io/parquet/experimental/variant_extract.cu | 8 +++++--- .../io/experimental/variant_extract_test.cpp | 15 ++++++++++----- 3 files changed, 15 insertions(+), 9 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant_spec.hpp b/cpp/include/cudf/io/experimental/variant_spec.hpp index b34da3564556..ed09a490fb43 100644 --- a/cpp/include/cudf/io/experimental/variant_spec.hpp +++ b/cpp/include/cudf/io/experimental/variant_spec.hpp @@ -48,7 +48,6 @@ enum class variant_primitive_type : uint8_t { /** * @brief Logical type of a VARIANT value as returned by get_variant_type_id. - * */ enum class variant_logical_type : uint8_t { OBJECT = 0, diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 24e5a7602ba6..a9921000fc51 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -792,9 +792,11 @@ struct cast_variant_fn { } }; -// Classifies only the first (value_metadata) byte of enc; does not validate the remaining payload. -// A recognized header returns its logical type even when the payload is truncated. -// Returns nullopt for an empty blob or an unrecognized primitive type ID. +/** + * @brief Classifies only the first (value_metadata) byte of enc; does not validate the remaining + * payload. A recognized header returns its logical type even when the payload is truncated. Returns + * nullopt for an empty blob or an unrecognized primitive type ID. + */ __device__ cuda::std::optional logical_type_of(device_span enc) { if (enc.empty()) { return cuda::std::nullopt; } diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 0b1b8ff82c73..d866589bdd9d 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1383,10 +1383,10 @@ TEST_F(InvalidInputShapeTest, CastVariantRejectsMalformedInput) } } -// get_variant_type_id requires a list input; every other shape must be rejected with -// std::invalid_argument. TEST_F(InvalidInputShapeTest, GetVariantTypeIdRejectsMalformedInput) { + // get_variant_type_id requires a list input; every other shape must be rejected with + // std::invalid_argument. auto stream = cudf::test::get_default_stream(); std::vector cases; @@ -1403,7 +1403,9 @@ TEST_F(InvalidInputShapeTest, GetVariantTypeIdRejectsMalformedInput) namespace { -// Helper: run get_variant_type_id on the value child of an apache fixture. +/** + * @brief Helper: run get_variant_type_id on the value child of an apache fixture. + */ template [[nodiscard]] std::unique_ptr apache_type_id(avf::fixture const& fixture) { @@ -1413,8 +1415,11 @@ template return cudf::io::parquet::experimental::get_variant_type_id(value, stream); } -// Build a list column from blobs with per-row validity. Rows where valid[i] is false are -// null at the list level (not an encoded Variant null — those are valid rows with a NULLVAL blob). +/** + * @brief Build a list column from blobs with per-row validity. Rows where valid[i] is false + * are null at the list level (not an encoded Variant null — those are valid rows with a NULLVAL + * blob). + */ [[nodiscard]] std::unique_ptr make_list_u8_nullable( cudf::host_span const> blobs, std::vector const& valid) { From b6e6d8fa8f71ed28c8c686af3e129c2a5eb6642b Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 11 Aug 2026 01:10:09 +0000 Subject: [PATCH 16/16] review --- cpp/include/cudf/io/experimental/variant.hpp | 2 +- .../parquet/experimental/variant_extract.cu | 12 +-- .../io/experimental/variant_extract_test.cpp | 98 ++++++++++--------- 3 files changed, 57 insertions(+), 55 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 53729920efa7..dca21e4c7da1 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -121,7 +121,7 @@ namespace io::parquet::experimental { * @param values `list` column of VARIANT-encoded value bytes * @param stream CUDA stream * @param mr Device memory resource - * @return `INT32` column of `variant_logical_type` values cast to `int32_t` + * @return `UINT8` column of `variant_logical_type` values cast to `uint8_t` * * @throws std::invalid_argument if `values` is not a `list` column */ diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index c631d8403290..f380538696ef 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -1005,7 +1005,7 @@ std::unique_ptr get_variant_type_id(column_view const& values, { validate_variant_child(values); size_type const num_rows = values.size(); - if (num_rows == 0) { return make_empty_column(data_type{type_id::INT32}); } + if (num_rows == 0) { return make_empty_column(data_type{type_id::UINT8}); } auto val_device_view = column_device_view::create(values, stream); cudf::lists_column_device_view val_lists_device_view(*val_device_view); @@ -1015,23 +1015,23 @@ std::unique_ptr get_variant_type_id(column_view const& values, : cudf::create_null_mask(num_rows, mask_state::ALL_VALID, stream, mr); auto* d_null_mask = static_cast(null_mask.data()); - rmm::device_buffer data{static_cast(num_rows) * sizeof(int32_t), stream, mr}; + rmm::device_buffer data{static_cast(num_rows) * sizeof(uint8_t), stream, mr}; thrust::transform( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::counting_iterator(0), cuda::counting_iterator(num_rows), - static_cast(data.data()), - [values = val_lists_device_view, d_null_mask] __device__(size_type row) -> int32_t { + static_cast(data.data()), + [values = val_lists_device_view, d_null_mask] __device__(size_type row) -> uint8_t { if (!cudf::bit_is_set(d_null_mask, row)) { return 0; } auto const ltype = logical_type_of(list_row_span(values, row)); - if (ltype.has_value()) { return static_cast(ltype.value()); } + if (ltype.has_value()) { return static_cast(ltype.value()); } cudf::clear_bit(d_null_mask, row); return 0; }); auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); - return std::make_unique(data_type{type_id::INT32}, + return std::make_unique(data_type{type_id::UINT8}, num_rows, std::move(data), null_count > 0 ? std::move(null_mask) : rmm::device_buffer{}, diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 7e20b51e3835..b6dd4d515549 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1463,13 +1463,13 @@ using LT = cudf::io::parquet::experimental::variant_logical_type; TEST_F(GetVariantTypeIdTest, NullValue) { auto got = apache_type_id(avf::primitive_null); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::NULL_VALUE)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::NULL_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, Boolean) { - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::BOOLEAN)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::BOOLEAN)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_boolean_true), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_boolean_false), expected); } @@ -1477,8 +1477,8 @@ TEST_F(GetVariantTypeIdTest, Boolean) TEST_F(GetVariantTypeIdTest, LongValueAllIntWidths) { // INT8, INT16, INT32, INT64 all map to long_value regardless of physical width. - cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::LONG_VALUE)}; + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::LONG_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int8), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int16), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_int32), expected); @@ -1488,7 +1488,7 @@ TEST_F(GetVariantTypeIdTest, LongValueAllIntWidths) TEST_F(GetVariantTypeIdTest, StringBothEncodings) { // SHORT_STRING and primitive LONG_STRING both map to string. - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::STRING)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::STRING)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::short_string), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_string), expected); } @@ -1497,20 +1497,20 @@ TEST_F(GetVariantTypeIdTest, FloatTypes) { { auto got = apache_type_id(avf::primitive_float); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::FLOAT_VALUE)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::FLOAT_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } { auto got = apache_type_id(avf::primitive_double); - cudf::test::fixed_width_column_wrapper expected{ - static_cast(LT::DOUBLE_VALUE)}; + cudf::test::fixed_width_column_wrapper expected{ + static_cast(LT::DOUBLE_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } } TEST_F(GetVariantTypeIdTest, Decimal) { - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::DECIMAL)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::DECIMAL)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal4), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal8), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_decimal16), expected); @@ -1519,15 +1519,15 @@ TEST_F(GetVariantTypeIdTest, Decimal) TEST_F(GetVariantTypeIdTest, Date) { auto got = apache_type_id(avf::primitive_date); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::DATE)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::DATE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, TimestampBothNanos) { // TIMESTAMP_MICROS and TIMESTAMP_NANOS both map to timestamp. - cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::TIMESTAMP)}; + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::TIMESTAMP)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestamp), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestamp_nanos), expected); } @@ -1535,8 +1535,8 @@ TEST_F(GetVariantTypeIdTest, TimestampBothNanos) TEST_F(GetVariantTypeIdTest, TimestampNtzBothNanos) { // TIMESTAMP_NTZ_MICROS and TIMESTAMP_NTZ_NANOS both map to timestamp_ntz. - cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::TIMESTAMP_NTZ)}; + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::TIMESTAMP_NTZ)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestampntz), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::primitive_timestampntz_nanos), expected); } @@ -1544,35 +1544,35 @@ TEST_F(GetVariantTypeIdTest, TimestampNtzBothNanos) TEST_F(GetVariantTypeIdTest, Binary) { auto got = apache_type_id(avf::primitive_binary); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::BINARY)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::BINARY)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, Uuid) { auto got = apache_type_id(avf::primitive_uuid); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::UUID)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::UUID)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, TimeNtz) { auto got = apache_type_id(avf::primitive_time); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::TIME_NTZ)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::TIME_NTZ)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } TEST_F(GetVariantTypeIdTest, ObjectAndArray) { { - cudf::test::fixed_width_column_wrapper const expected{ - static_cast(LT::OBJECT)}; + cudf::test::fixed_width_column_wrapper const expected{ + static_cast(LT::OBJECT)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_primitive), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_nested), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::object_empty), expected); } { - cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::ARRAY)}; + cudf::test::fixed_width_column_wrapper const expected{static_cast(LT::ARRAY)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_primitive), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_nested), expected); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*apache_type_id(avf::array_empty), expected); @@ -1603,9 +1603,10 @@ TEST_F(GetVariantTypeIdTest, InputNullRowPropagates) auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); - cudf::test::fixed_width_column_wrapper expected( - {static_cast(LT::LONG_VALUE), 0, static_cast(LT::LONG_VALUE)}, - {true, false, true}); + std::initializer_list expected_vals1{ + static_cast(LT::LONG_VALUE), 0, static_cast(LT::LONG_VALUE)}; + std::initializer_list validity1{true, false, true}; + cudf::test::fixed_width_column_wrapper expected(expected_vals1, validity1); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1619,7 +1620,7 @@ TEST_F(GetVariantTypeIdTest, EncodedNullIsNotInputNull) ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 0); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::NULL_VALUE)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::NULL_VALUE)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1632,9 +1633,10 @@ TEST_F(GetVariantTypeIdTest, EmptyValueBlobProducesNull) auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); - cudf::test::fixed_width_column_wrapper expected( - {static_cast(LT::LONG_VALUE), 0, static_cast(LT::LONG_VALUE)}, - {true, false, true}); + std::initializer_list expected_vals2{ + static_cast(LT::LONG_VALUE), 0, static_cast(LT::LONG_VALUE)}; + std::initializer_list validity2{true, false, true}; + cudf::test::fixed_width_column_wrapper expected(expected_vals2, validity2); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1661,12 +1663,12 @@ TEST_F(GetVariantTypeIdTest, MixedTypesColumn) }; auto got = cudf::io::parquet::experimental::get_variant_type_id(values, stream); - cudf::test::fixed_width_column_wrapper expected{ - static_cast(LT::NULL_VALUE), - static_cast(LT::BOOLEAN), - static_cast(LT::LONG_VALUE), - static_cast(LT::STRING), - static_cast(LT::DOUBLE_VALUE), + cudf::test::fixed_width_column_wrapper expected{ + static_cast(LT::NULL_VALUE), + static_cast(LT::BOOLEAN), + static_cast(LT::LONG_VALUE), + static_cast(LT::STRING), + static_cast(LT::DOUBLE_VALUE), }; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1692,7 +1694,7 @@ TEST_F(GetVariantTypeIdTest, EmptyInput) cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); - EXPECT_EQ(got->type().id(), cudf::type_id::INT32); + EXPECT_EQ(got->type().id(), cudf::type_id::UINT8); EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); } @@ -1708,8 +1710,8 @@ TEST_F(GetVariantTypeIdTest, SlicedValuesColumn) auto const sliced_values = cudf::slice(value_child, {1, 3}).front(); auto got = cudf::io::parquet::experimental::get_variant_type_id(sliced_values, stream); - cudf::test::fixed_width_column_wrapper expected{static_cast(LT::OBJECT), - static_cast(LT::OBJECT)}; + cudf::test::fixed_width_column_wrapper expected{static_cast(LT::OBJECT), + static_cast(LT::OBJECT)}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1721,23 +1723,23 @@ TEST_F(GetVariantTypeIdTest, LargeMultiRowColumn) struct row_spec { std::vector blob; - int32_t expected_id; + uint8_t expected_id; }; std::vector const types{ - {enc_null(), static_cast(LT::NULL_VALUE)}, - {enc_bool(false), static_cast(LT::BOOLEAN)}, - {enc_int8(1), static_cast(LT::LONG_VALUE)}, - {enc_int16(2), static_cast(LT::LONG_VALUE)}, - {enc_int32(3), static_cast(LT::LONG_VALUE)}, - {enc_int64(4), static_cast(LT::LONG_VALUE)}, - {enc_float64(5.0), static_cast(LT::DOUBLE_VALUE)}, - {enc_short_string("x"), static_cast(LT::STRING)}, - {enc_long_string(std::string(70, 'z')), static_cast(LT::STRING)}, + {enc_null(), static_cast(LT::NULL_VALUE)}, + {enc_bool(false), static_cast(LT::BOOLEAN)}, + {enc_int8(1), static_cast(LT::LONG_VALUE)}, + {enc_int16(2), static_cast(LT::LONG_VALUE)}, + {enc_int32(3), static_cast(LT::LONG_VALUE)}, + {enc_int64(4), static_cast(LT::LONG_VALUE)}, + {enc_float64(5.0), static_cast(LT::DOUBLE_VALUE)}, + {enc_short_string("x"), static_cast(LT::STRING)}, + {enc_long_string(std::string(70, 'z')), static_cast(LT::STRING)}, }; constexpr int num_rows = 600; std::vector> blobs(num_rows); - std::vector expected_ids(num_rows); + std::vector expected_ids(num_rows); for (int i = 0; i < num_rows; ++i) { auto const& spec = types[i % types.size()]; blobs[i] = spec.blob; @@ -1747,7 +1749,7 @@ TEST_F(GetVariantTypeIdTest, LargeMultiRowColumn) auto values = make_list_u8_nullable(blobs, std::vector(num_rows, true)); auto got = cudf::io::parquet::experimental::get_variant_type_id(*values, stream); - cudf::test::fixed_width_column_wrapper expected(expected_ids.begin(), + cudf::test::fixed_width_column_wrapper expected(expected_ids.begin(), expected_ids.end()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); }