From 64ea90970dd997960f4db98826b016b19ddd50f8 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Wed, 15 Jul 2026 19:13:24 +0000 Subject: [PATCH 1/9] first commit --- cpp/include/cudf/io/experimental/variant.hpp | 6 +- .../parquet/experimental/variant_extract.cu | 78 ++++++++++++++++++- .../io/experimental/variant_extract_test.cpp | 20 ++++- 3 files changed, 97 insertions(+), 7 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 4ee7747dc055..f67da9cef2e0 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -71,13 +71,13 @@ namespace io::parquet::experimental { * `desired_type`. * * @param values `list` column of VARIANT-encoded value bytes - * @param desired_type Target cuDF type (`STRING` or `INT8`/`INT16`/`INT32`/`INT64`) + * @param desired_type Target cuDF type (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or `BOOL8`) * @param stream CUDA stream * @param mr Device memory resource * @return Typed column decoded from the VARIANT value blobs * * @throws std::invalid_argument if `values` is not a `list` column, or if `desired_type` - * is not one of the supported types (`STRING` or `INT8`/`INT16`/`INT32`/`INT64`) + * is not one of the supported types (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or `BOOL8`) */ [[nodiscard]] std::unique_ptr cast_variant( column_view const& values, @@ -93,7 +93,7 @@ namespace io::parquet::experimental { * * @param variant_column Struct column (VARIANT materialization) * @param path JSONPath-like path string (see `get_variant_field` for syntax) - * @param desired_type Target type: `STRING` or `INT8`/`INT16`/`INT32`/`INT64` + * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or `BOOL8` * @param stream CUDA stream * @param mr Device memory resource * @return Column of `desired_type` diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index aabc55449405..83f79a2b0619 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -376,10 +376,11 @@ constexpr bool is_variant_int = cuda::std::is_same_v || cuda::std::is_same_v || cuda::std::is_same_v || cuda::std::is_same_v; -// The output types a VARIANT value can be cast to: the fixed-width signed integers plus strings. +// The output types a VARIANT value can be cast to: the fixed-width signed integers, bool, and +// strings. template constexpr bool is_variant_castable = - is_variant_int || cuda::std::is_same_v; + is_variant_int || cuda::std::is_same_v || cuda::std::is_same_v; // Variant primitive ints: basic_type == primitive, value_header maps INT{8,16,32,64}. template @@ -401,6 +402,23 @@ __device__ inline cuda::std::optional decode_int(device_span e return cudf::io::unaligned_load(enc.data() + 1); } +/** + * @brief Decode a single VARIANT value blob into a bool. + * + * Boolean values carry no payload: the distinction between true and false is encoded entirely in + * the primitive type header (`boolean_true` vs `boolean_false`). + */ +__device__ inline cuda::std::optional decode_bool(device_span enc) +{ + if (enc.size() < 1) { return cuda::std::nullopt; } + uint8_t const value_metadata = enc[0]; + if (variant_basic_type(value_metadata) != basic_type::primitive) { return cuda::std::nullopt; } + auto const value_header = variant_value_header(value_metadata); + if (value_header == static_cast(primitive_type::boolean_true)) { return true; } + if (value_header == static_cast(primitive_type::boolean_false)) { return false; } + return cuda::std::nullopt; +} + __device__ device_span resolve_path(device_span meta, device_span val, column_device_view path) @@ -543,6 +561,42 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_int_kernel( } } +/** + * @brief Per-row kernel: decode each VARIANT value blob into a bool. + * + * Boolean values are encoded as a single-byte `primitive_type::boolean_true` or + * `primitive_type::boolean_false` header with no payload. Rows that are null, or whose value is + * not a boolean primitive, are marked null in `d_null_mask` with an output of false. + */ +CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_bool_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] = false; + continue; + } + + auto const val_begin = values.offset_at(row); + auto const val_end = values.offset_at(row + 1); + auto const val_child = values.child(); + device_span const val{val_child.data() + val_begin, + static_cast(val_end - val_begin)}; + + auto const decoded = decode_bool(val); + if (decoded.has_value()) { + d_output[row] = *decoded; + } else { + d_output[row] = false; + cudf::clear_bit(d_null_mask, row); + } + } +} + /** * @brief Strings-children functor: decode each VARIANT value blob into a string. * @@ -625,6 +679,26 @@ struct cast_variant_fn { null_count); } + template + std::unique_ptr operator()() + requires(cuda::std::is_same_v) + { + rmm::device_buffer data{num_rows * sizeof(bool), stream, mr}; + + auto grid = cudf::detail::grid_1d{num_rows, block_size}; + cast_variant_bool_kernel<<>>( + values, {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(desired_type, + num_rows, + std::move(data), + null_count > 0 ? std::move(null_mask) : rmm::device_buffer{}, + null_count); + } + template std::unique_ptr operator()() requires(cuda::std::is_same_v) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 051c3d94e240..1883a3102e7a 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -658,6 +658,22 @@ TEST_F(CastVariantTest, ApachePrimitiveInts) cast(avf::primitive_int64, int64_t{1234567890123456789LL}); } +TEST_F(CastVariantTest, ApachePrimitiveBooleans) +{ + auto stream = cudf::test::get_default_stream(); + auto const cast = [&](auto const& fixture, bool expected_val) { + auto col = make_apache_variant(fixture); + auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); + auto got = cudf::io::parquet::experimental::cast_variant( + value, cudf::data_type{cudf::type_id::BOOL8}, stream); + cudf::test::fixed_width_column_wrapper expected{expected_val}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + }; + + cast(avf::primitive_boolean_true, true); + cast(avf::primitive_boolean_false, false); +} + TEST_F(CastVariantTest, ApacheShortString) { auto col = make_apache_variant(avf::short_string); @@ -709,7 +725,7 @@ TEST_F(CastVariantTest, EmptyInput) auto const values = cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); - for (auto const id : {cudf::type_id::INT32, cudf::type_id::STRING}) { + for (auto const id : {cudf::type_id::INT32, cudf::type_id::STRING, cudf::type_id::BOOL8}) { auto got = cudf::io::parquet::experimental::cast_variant(*values, cudf::data_type{id}, stream); EXPECT_EQ(got->type().id(), id); EXPECT_EQ(got->size(), 0); From 1570390035904fe94a93551e7368c4fbef8377ca Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Fri, 24 Jul 2026 17:23:22 +0000 Subject: [PATCH 2/9] rebasing bools on top of floats --- cpp/include/cudf/io/experimental/variant.hpp | 9 ++- .../parquet/experimental/variant_extract.cu | 71 +++++++++++++------ cpp/src/io/utilities/block_utils.cuh | 4 +- .../io/experimental/variant_extract_test.cpp | 23 +++++- 4 files changed, 78 insertions(+), 29 deletions(-) diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index f67da9cef2e0..bba41b22e16d 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -71,13 +71,15 @@ namespace io::parquet::experimental { * `desired_type`. * * @param values `list` column of VARIANT-encoded value bytes - * @param desired_type Target cuDF type (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or `BOOL8`) + * @param desired_type Target cuDF type (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, + * `FLOAT32`/`FLOAT64`, or `BOOL8`) * @param stream CUDA stream * @param mr Device memory resource * @return Typed column decoded from the VARIANT value blobs * * @throws std::invalid_argument if `values` is not a `list` column, or if `desired_type` - * is not one of the supported types (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or `BOOL8`) + * is not one of the supported types (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, + * `FLOAT32`/`FLOAT64`, or `BOOL8`) */ [[nodiscard]] std::unique_ptr cast_variant( column_view const& values, @@ -93,7 +95,8 @@ namespace io::parquet::experimental { * * @param variant_column Struct column (VARIANT materialization) * @param path JSONPath-like path string (see `get_variant_field` for syntax) - * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or `BOOL8` + * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, + * `FLOAT32`/`FLOAT64`, or `BOOL8` * @param stream CUDA stream * @param mr Device memory resource * @return Column of `desired_type` diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 83f79a2b0619..8c42388e3445 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -373,30 +373,54 @@ __device__ device_span locate_object_field(device_span constexpr bool is_variant_int = - cuda::std::is_same_v || cuda::std::is_same_v || - cuda::std::is_same_v || cuda::std::is_same_v; + cudf::is_integral_not_bool() && cudf::is_signed() && !cuda::std::is_same_v; -// The output types a VARIANT value can be cast to: the fixed-width signed integers, bool, and -// strings. +// The fixed-width primitive types (signed integers and floats) a VARIANT value can be decoded into. template -constexpr bool is_variant_castable = - is_variant_int || cuda::std::is_same_v || cuda::std::is_same_v; +constexpr bool is_variant_primitive = is_variant_int || cudf::is_floating_point(); -// Variant primitive ints: basic_type == primitive, value_header maps INT{8,16,32,64}. +// The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, bool, +// and strings. template -__device__ inline cuda::std::optional decode_int(device_span enc) +constexpr bool is_variant_castable = is_variant_primitive || cuda::std::is_same_v || + cuda::std::is_same_v; + +// Maps a fixed-width output type to the VARIANT primitive type header id that encodes it. +template + requires(is_variant_primitive) +__device__ constexpr primitive_type primitive_type_for() { - static_assert(is_variant_int, "decode_int: T must be int8_t, int16_t, int32_t, or int64_t"); + if constexpr (cuda::std::is_same_v) { + return primitive_type::int8; + } else if constexpr (cuda::std::is_same_v) { + return primitive_type::int16; + } else if constexpr (cuda::std::is_same_v) { + return primitive_type::int32; + } else if constexpr (cuda::std::is_same_v) { + return primitive_type::int64; + } else if constexpr (cuda::std::is_same_v) { + return primitive_type::float32; + } else if constexpr (cuda::std::is_same_v) { + return primitive_type::float64; + } else { + CUDF_UNREACHABLE("primitive_type_for: T is not a supported variant primitive type"); + return primitive_type::null; + } +} +/** + * @brief Decode a single VARIANT value blob into a fixed-width primitive of type `T`. + * + * Requires `basic_type == primitive` and a value header whose physical type id matches `T` exactly. + */ +template +__device__ inline cuda::std::optional decode_primitive(device_span enc) +{ if (cuda::std::cmp_less(enc.size(), 1 + sizeof(T))) { return cuda::std::nullopt; } - constexpr primitive_type expected = cuda::std::is_same_v ? primitive_type::int8 - : cuda::std::is_same_v ? primitive_type::int16 - : cuda::std::is_same_v ? primitive_type::int32 - : primitive_type::int64; - uint8_t const value_metadata = enc[0]; + uint8_t const value_metadata = enc[0]; if (variant_basic_type(value_metadata) != basic_type::primitive || - variant_value_header(value_metadata) != static_cast(expected)) { + variant_value_header(value_metadata) != static_cast(primitive_type_for())) { return cuda::std::nullopt; } return cudf::io::unaligned_load(enc.data() + 1); @@ -524,15 +548,16 @@ CUDF_KERNEL __launch_bounds__(block_size) void locate_variant_fields_kernel( } /** - * @brief Per-row kernel: decode each VARIANT value blob into an integer of type `T`. + * @brief Per-row kernel: decode each VARIANT value blob into a fixed-width primitive of type `T`. * * Writes the decoded value to `d_output[row]` for non-null rows whose blob is a variant primitive - * int whose physical type id matches `T` exactly (e.g. an int16 value does not decode into an - * int32 output; there is no widening). Rows that are null, or whose value is not an exact-width - * match for `T`, are marked null in `d_null_mask` with an output of 0. + * whose physical type id matches `T` exactly (e.g. an int16 value does not decode into an int32 + * output, and a float32 value does not decode into a float64 output; there is no widening). Rows + * that are null, or whose value is not an exact-width match for `T`, are marked null in + * `d_null_mask` with an output of 0. */ template -CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_int_kernel( +CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_primitive_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()); @@ -551,7 +576,7 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_int_kernel( device_span const val{val_child.data() + val_begin, static_cast(val_end - val_begin)}; - auto const decoded = decode_int(val); + auto const decoded = decode_primitive(val); if (decoded.has_value()) { d_output[row] = *decoded; } else { @@ -661,12 +686,12 @@ struct cast_variant_fn { template std::unique_ptr operator()() - requires(is_variant_int) + requires(is_variant_primitive) { rmm::device_buffer data{num_rows * sizeof(T), stream, mr}; auto grid = cudf::detail::grid_1d{num_rows, block_size}; - cast_variant_int_kernel<<>>( + cast_variant_primitive_kernel<<>>( values, {static_cast(data.data()), static_cast(num_rows)}, d_null_mask); CUDF_CUDA_TRY(cudaGetLastError()); diff --git a/cpp/src/io/utilities/block_utils.cuh b/cpp/src/io/utilities/block_utils.cuh index d00e91ef2bb0..d157378127d3 100644 --- a/cpp/src/io/utilities/block_utils.cuh +++ b/cpp/src/io/utilities/block_utils.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -58,7 +58,7 @@ inline __device__ T warp_reduce_pos(T pos, uint32_t t) } template - requires(cuda::std::is_integral_v) + requires(cuda::std::is_trivially_copyable_v) inline __device__ T unaligned_load(uint8_t const* p) { T value; diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 1883a3102e7a..7bff24d79da5 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -658,6 +658,23 @@ TEST_F(CastVariantTest, ApachePrimitiveInts) cast(avf::primitive_int64, int64_t{1234567890123456789LL}); } +TEST_F(CastVariantTest, ApachePrimitiveFloats) +{ + auto stream = cudf::test::get_default_stream(); + auto const cast = [&](auto const& fixture, auto expected_val) { + using T = decltype(expected_val); + auto col = make_apache_variant(fixture); + auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); + auto got = cudf::io::parquet::experimental::cast_variant( + value, cudf::data_type{cudf::type_to_id()}, stream); + cudf::test::fixed_width_column_wrapper expected{expected_val}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + }; + + cast(avf::primitive_float, float{1234567936.0f}); + cast(avf::primitive_double, double{1234567890.1234}); +} + TEST_F(CastVariantTest, ApachePrimitiveBooleans) { auto stream = cudf::test::get_default_stream(); @@ -725,7 +742,11 @@ TEST_F(CastVariantTest, EmptyInput) auto const values = cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); - for (auto const id : {cudf::type_id::INT32, cudf::type_id::STRING, cudf::type_id::BOOL8}) { + for (auto const id : {cudf::type_id::INT32, + cudf::type_id::STRING, + cudf::type_id::FLOAT32, + cudf::type_id::FLOAT64, + cudf::type_id::BOOL8}) { auto got = cudf::io::parquet::experimental::cast_variant(*values, cudf::data_type{id}, stream); EXPECT_EQ(got->type().id(), id); EXPECT_EQ(got->size(), 0); From e87a6e870f50b71bbb2597c604bc33d412b6a030 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 28 Jul 2026 20:56:57 +0000 Subject: [PATCH 3/9] adds boolean tests --- .../io/experimental/variant_extract_test.cpp | 58 +++++++++++++++++++ 1 file changed, 58 insertions(+) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index ae0d7765a63d..72e4e751edee 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -839,6 +839,64 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) cast(avf::primitive_boolean_true, true); cast(avf::primitive_boolean_false, false); + + // Null variant value must cast to a null BOOL8, not false. + { + auto col = make_apache_variant(avf::primitive_null); + auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); + auto got = cudf::io::parquet::experimental::cast_variant( + value, cudf::data_type{cudf::type_id::BOOL8}, stream); + cudf::test::fixed_width_column_wrapper expected({false}, {false}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } + + // Sliced multi-row column exercises grid-stride paths with a non-zero slice offset. + { + constexpr int num_rows = 130; + constexpr int slice_beg = 2; + constexpr int slice_end = 128; + + std::vector const true_bytes{avf::primitive_boolean_true.value.begin(), + avf::primitive_boolean_true.value.end()}; + std::vector const false_bytes{avf::primitive_boolean_false.value.begin(), + avf::primitive_boolean_false.value.end()}; + std::vector const null_bytes{avf::primitive_null.value.begin(), + avf::primitive_null.value.end()}; + std::vector const meta_bytes{avf::primitive_boolean_true.metadata.begin(), + avf::primitive_boolean_true.metadata.end()}; + + std::vector> metas(num_rows, meta_bytes); + std::vector> vals(num_rows); + std::vector exp_vals(num_rows); + std::vector exp_valid(num_rows); + + for (int i = 0; i < num_rows; ++i) { + int const pat = i % 3; + if (pat == 0) { + vals[i] = true_bytes; + exp_vals[i] = true; + exp_valid[i] = true; + } else if (pat == 1) { + vals[i] = false_bytes; + exp_vals[i] = false; + exp_valid[i] = true; + } else { + vals[i] = null_bytes; + exp_vals[i] = false; + exp_valid[i] = false; + } + } + + auto col = wrap_multi_row_variant(metas, vals); + auto const sliced = cudf::slice(col, {slice_beg, slice_end}).front(); + auto const value = cudf::structs_column_view{sliced}.get_sliced_child(1, stream); + auto got = cudf::io::parquet::experimental::cast_variant( + value, cudf::data_type{cudf::type_id::BOOL8}, stream); + + cudf::test::fixed_width_column_wrapper expected( + exp_vals.begin() + slice_beg, exp_vals.begin() + slice_end, exp_valid.begin() + slice_beg); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); + } } TEST_F(CastVariantTest, ApacheShortString) From 0bab85c40b8257f0e8f0024d3d39436ac557272f Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 30 Jul 2026 17:10:41 -0500 Subject: [PATCH 4/9] Update cpp/src/io/parquet/experimental/variant_extract.cu Co-authored-by: Vukasin Milovanovic --- cpp/src/io/parquet/experimental/variant_extract.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 696b227f71ac..ff93d351d138 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -472,7 +472,7 @@ __device__ inline cuda::std::optional decode_primitive(device_span decode_bool(device_span enc) { - if (enc.size() < 1) { return cuda::std::nullopt; } + if (enc.empty()) { return cuda::std::nullopt; } uint8_t const value_metadata = enc[0]; if (variant_basic_type(value_metadata) != basic_type::primitive) { return cuda::std::nullopt; } auto const value_header = variant_value_header(value_metadata); From 6e7ba18e1e12478db40c4373da223f3bec456cc6 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 30 Jul 2026 22:11:10 +0000 Subject: [PATCH 5/9] changing primitive --- cpp/src/io/parquet/experimental/variant_extract.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index e619ade55d9a..84c8cd38d5e6 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -436,17 +436,17 @@ constexpr bool is_variant_int = // The fixed-width primitive types (signed integers and floats) a VARIANT value can be decoded into. template -constexpr bool is_variant_primitive = is_variant_int || cudf::is_floating_point(); +constexpr bool is_variant_numerical = is_variant_int || cudf::is_floating_point(); // The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, bool, // and strings. template -constexpr bool is_variant_castable = is_variant_primitive || cuda::std::is_same_v || +constexpr bool is_variant_castable = is_variant_numerical || cuda::std::is_same_v || cuda::std::is_same_v; // Maps a fixed-width output type to the VARIANT primitive type header id that encodes it. template - requires(is_variant_primitive) + requires(is_variant_numerical) __device__ constexpr primitive_type primitive_type_for() { if constexpr (cuda::std::is_same_v) { @@ -783,7 +783,7 @@ struct cast_variant_fn { template std::unique_ptr operator()() - requires(is_variant_primitive) + requires(is_variant_numerical) { rmm::device_buffer data{num_rows * sizeof(T), stream, mr}; From a1b1aeb16a19299e185823e66263de6a7dbb7673 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 30 Jul 2026 22:45:07 +0000 Subject: [PATCH 6/9] reviews --- .../parquet/experimental/variant_extract.cu | 41 +++++++------------ 1 file changed, 15 insertions(+), 26 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 4dba27d7e84a..766b9a2fa32c 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -474,10 +474,10 @@ __device__ inline cuda::std::optional decode_bool(device_span(primitive_type::boolean_true)) { return true; } - if (value_header == static_cast(primitive_type::boolean_false)) { return false; } + if (value_header == static_cast(primitive_type::BOOLEAN_TRUE)) { return true; } + if (value_header == static_cast(primitive_type::BOOLEAN_FALSE)) { return false; } return cuda::std::nullopt; } @@ -564,20 +564,21 @@ __device__ cuda::std::optional> decode_string( return cuda::std::nullopt; } +__device__ device_span list_row_span(cudf::lists_column_device_view const& col, + size_type row) +{ + auto const begin = col.offset_at(row); + auto const end = col.offset_at(row + 1); + return {col.child().data() + begin, static_cast(end - begin)}; +} + // Returns the metadata and value list bytes for a given row from device views __device__ cuda::std::pair, device_span> metadata_and_value_at(cudf::lists_column_device_view const& metadata, cudf::lists_column_device_view const& values, size_type row) { - auto const meta_begin = metadata.offset_at(row); - auto const meta_end = metadata.offset_at(row + 1); - auto const val_begin = values.offset_at(row); - auto const val_end = values.offset_at(row + 1); - return { - {metadata.child().data() + meta_begin, - static_cast(meta_end - meta_begin)}, - {values.child().data() + val_begin, static_cast(val_end - val_begin)}}; + return {list_row_span(metadata, row), list_row_span(values, row)}; } constexpr int block_size = 256; @@ -646,11 +647,7 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_primitive_kernel( continue; } - auto const val_begin = values.offset_at(row); - auto const val_end = values.offset_at(row + 1); - auto const val_child = values.child(); - device_span const val{val_child.data() + val_begin, - static_cast(val_end - val_begin)}; + auto const val = list_row_span(values, row); auto const decoded = decode_primitive(val); if (decoded.has_value()) { @@ -682,11 +679,7 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_bool_kernel( continue; } - auto const val_begin = values.offset_at(row); - auto const val_end = values.offset_at(row + 1); - auto const val_child = values.child(); - device_span const val{val_child.data() + val_begin, - static_cast(val_end - val_begin)}; + auto const val = list_row_span(values, row); auto const decoded = decode_bool(val); if (decoded.has_value()) { @@ -720,11 +713,7 @@ struct cast_variant_string_fn { return; } - auto const val_begin = d_values.offset_at(row); - auto const val_end = d_values.offset_at(row + 1); - auto const val_child = d_values.child(); - device_span const val{val_child.data() + val_begin, - static_cast(val_end - val_begin)}; + auto const val = list_row_span(d_values, row); auto const str = decode_string(val); if (!str) { From 69bd1d08a47074fd4fab6b424ea1df4c961ebb2d Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Thu, 30 Jul 2026 23:07:02 +0000 Subject: [PATCH 7/9] reviews --- cpp/tests/io/experimental/variant_extract_test.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 7ab4f53cca39..e810bbe142a0 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -999,10 +999,12 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) } // Sliced multi-row column exercises grid-stride paths with a non-zero slice offset. + // num_rows / slice range is chosen so the sliced window (slice_end - slice_beg = 512) + // spans more than one cast_variant_bool_kernel block (block_size = 256). { - constexpr int num_rows = 130; - constexpr int slice_beg = 2; - constexpr int slice_end = 128; + constexpr int num_rows = 516; + constexpr int slice_beg = 3; + constexpr int slice_end = 515; std::vector const true_bytes{avf::primitive_boolean_true.value.begin(), avf::primitive_boolean_true.value.end()}; From 2ad01811fcef2125905345e9c0724a61593f7437 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 4 Aug 2026 19:10:02 +0000 Subject: [PATCH 8/9] addressing reviews --- .../parquet/experimental/variant_extract.cu | 51 ++++++------------- 1 file changed, 15 insertions(+), 36 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 766b9a2fa32c..3c19bd80295b 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -40,6 +40,8 @@ #include #include #include +#include +#include #include #include @@ -659,38 +661,6 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_primitive_kernel( } } -/** - * @brief Per-row kernel: decode each VARIANT value blob into a bool. - * - * Boolean values are encoded as a single-byte `primitive_type::boolean_true` or - * `primitive_type::boolean_false` header with no payload. Rows that are null, or whose value is - * not a boolean primitive, are marked null in `d_null_mask` with an output of false. - */ -CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_bool_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] = false; - continue; - } - - auto const val = list_row_span(values, row); - - auto const decoded = decode_bool(val); - if (decoded.has_value()) { - d_output[row] = *decoded; - } else { - d_output[row] = false; - cudf::clear_bit(d_null_mask, row); - } - } -} - /** * @brief Strings-children functor: decode each VARIANT value blob into a string. * @@ -775,10 +745,19 @@ struct cast_variant_fn { { rmm::device_buffer data{num_rows * sizeof(bool), stream, mr}; - auto grid = cudf::detail::grid_1d{num_rows, block_size}; - cast_variant_bool_kernel<<>>( - values, {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()), + thrust::counting_iterator(0), + thrust::counting_iterator(num_rows), + static_cast(data.data()), + [values = this->values, d_null_mask = this->d_null_mask] __device__(size_type row) -> bool { + if (!cudf::bit_is_set(d_null_mask, row)) { return false; } + auto const val = list_row_span(values, row); + auto const decoded = decode_bool(val); + if (decoded.has_value()) { return *decoded; } + cudf::clear_bit(d_null_mask, row); + return false; + }); auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); From ddf1b6895ca2b394b69e5602711feeede3ac4899 Mon Sep 17 00:00:00 2001 From: Abigale Kim Date: Tue, 4 Aug 2026 23:52:27 +0000 Subject: [PATCH 9/9] reviews --- cpp/src/io/parquet/experimental/variant_extract.cu | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 3c19bd80295b..f463230489b5 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -34,14 +34,13 @@ #include #include +#include #include #include #include #include #include #include -#include -#include #include #include @@ -747,8 +746,8 @@ struct cast_variant_fn { thrust::transform( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - thrust::counting_iterator(0), - thrust::counting_iterator(num_rows), + cuda::counting_iterator(0), + cuda::counting_iterator(num_rows), static_cast(data.data()), [values = this->values, d_null_mask = this->d_null_mask] __device__(size_type row) -> bool { if (!cudf::bit_is_set(d_null_mask, row)) { return false; }