From 6ff24c5aff2d8fd7cdb0898f31d0bb1892c0feba Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Wed, 1 Jul 2026 22:22:53 +0000
Subject: [PATCH 01/11] float pr

---
 cpp/include/cudf/io/experimental/variant.hpp  | 11 ++-
 .../parquet/experimental/variant_extract.cu   | 94 ++++++++++++++++++-
 .../io/experimental/variant_extract_test.cpp  | 24 ++++-
 3 files changed, 121 insertions(+), 8 deletions(-)

diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp
index 62763c8588d4..857c26cd82e9 100644
--- a/cpp/include/cudf/io/experimental/variant.hpp
+++ b/cpp/include/cudf/io/experimental/variant.hpp
@@ -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
  */
 
@@ -68,13 +68,15 @@ namespace io::parquet::experimental {
  * `desired_type`.
  *
  * @param values `list<uint8>` 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
+ *        `FLOAT32`/`FLOAT64`)
  * @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<uint8>` 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
+ *         `FLOAT32`/`FLOAT64`)
  */
 [[nodiscard]] std::unique_ptr<column> cast_variant(
   column_view const& values,
@@ -90,7 +92,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` or `INT8`/`INT16`/`INT32`/`INT64`
+ * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, or
+ *        `FLOAT32`/`FLOAT64`
  * @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..48352a723666 100644
--- a/cpp/src/io/parquet/experimental/variant_extract.cu
+++ b/cpp/src/io/parquet/experimental/variant_extract.cu
@@ -376,10 +376,15 @@ constexpr bool is_variant_int =
   cuda::std::is_same_v<T, int8_t> || cuda::std::is_same_v<T, int16_t> ||
   cuda::std::is_same_v<T, int32_t> || cuda::std::is_same_v<T, int64_t>;
 
-// The output types a VARIANT value can be cast to: the fixed-width signed integers plus strings.
+// The floating-point types a VARIANT value can be cast to: FLOAT{32,64}.
+template <typename T>
+constexpr bool is_variant_float = cuda::std::is_same_v<T, float> || cuda::std::is_same_v<T, double>;
+
+// The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, and
+// strings.
 template <typename T>
 constexpr bool is_variant_castable =
-  is_variant_int<T> || cuda::std::is_same_v<T, cudf::string_view>;
+  is_variant_int<T> || is_variant_float<T> || cuda::std::is_same_v<T, cudf::string_view>;
 
 // Variant primitive ints: basic_type == primitive, value_header maps INT{8,16,32,64}.
 template <typename T>
@@ -401,6 +406,33 @@ __device__ inline cuda::std::optional<T> decode_int(device_span<uint8_t const> e
   return cudf::io::unaligned_load<T>(enc.data() + 1);
 }
 
+/**
+ * @brief Decode a single VARIANT value blob into a float of type `T`.
+ *
+ * Matches only when the blob is a primitive value whose physical type id is the exact-width float
+ * header for `T` (FLOAT32 for `float`, FLOAT64 for `double`); there is no widening between float32
+ * and float64. `unaligned_load` is constrained to integral types, so the payload is copied via
+ * `memcpy` instead.
+ */
+template <typename T>
+__device__ inline cuda::std::optional<T> decode_float(device_span<uint8_t const> enc)
+{
+  static_assert(is_variant_float<T>, "decode_float: T must be float or double");
+
+  if (cuda::std::cmp_less(enc.size(), 1 + sizeof(T))) { return cuda::std::nullopt; }
+
+  constexpr primitive_type expected =
+    cuda::std::is_same_v<T, float> ? primitive_type::float32 : primitive_type::float64;
+  uint8_t const value_metadata = enc[0];
+  if (variant_basic_type(value_metadata) != basic_type::primitive ||
+      variant_value_header(value_metadata) != static_cast<uint8_t>(expected)) {
+    return cuda::std::nullopt;
+  }
+  T value;
+  cuda::std::memcpy(&value, enc.data() + 1, sizeof(T));
+  return value;
+}
+
 __device__ device_span<uint8_t const> resolve_path(device_span<uint8_t const> meta,
                                                    device_span<uint8_t const> val,
                                                    column_device_view path)
@@ -543,6 +575,44 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_int_kernel(
   }
 }
 
+/**
+ * @brief Per-row kernel: decode each VARIANT value blob into a float of type `T`.
+ *
+ * Writes the decoded value to `d_output[row]` for non-null rows whose blob is a variant primitive
+ * float whose physical type id matches `T` exactly (e.g. 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 <typename T>
+CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_float_kernel(
+  cudf::lists_column_device_view values, device_span<T> d_output, bitmask_type* d_null_mask)
+{
+  auto const num_rows = static_cast<size_type>(d_output.size());
+  auto const tid      = cudf::detail::grid_1d::global_thread_id<block_size>();
+  auto const stride   = cudf::detail::grid_1d::grid_stride<block_size>();
+
+  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);
+    auto const val_child = values.child();
+    device_span<uint8_t const> const val{val_child.data<uint8_t>() + val_begin,
+                                         static_cast<std::size_t>(val_end - val_begin)};
+
+    auto const decoded = decode_float<T>(val);
+    if (decoded.has_value()) {
+      d_output[row] = *decoded;
+    } else {
+      d_output[row] = 0;
+      cudf::clear_bit(d_null_mask, row);
+    }
+  }
+}
+
 /**
  * @brief Strings-children functor: decode each VARIANT value blob into a string.
  *
@@ -625,6 +695,26 @@ struct cast_variant_fn {
                                     null_count);
   }
 
+  template <typename T>
+  std::unique_ptr<column> operator()()
+    requires(is_variant_float<T>)
+  {
+    rmm::device_buffer data{num_rows * sizeof(T), stream, mr};
+
+    auto grid = cudf::detail::grid_1d{num_rows, block_size};
+    cast_variant_float_kernel<T><<<grid.num_blocks, block_size, 0, stream.value()>>>(
+      values, {static_cast<T*>(data.data()), static_cast<std::size_t>(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<column>(desired_type,
+                                    num_rows,
+                                    std::move(data),
+                                    null_count > 0 ? std::move(null_mask) : rmm::device_buffer{},
+                                    null_count);
+  }
+
   template <typename T>
   std::unique_ptr<column> operator()()
     requires(cuda::std::is_same_v<T, cudf::string_view>)
diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp
index 051c3d94e240..157605272878 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,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<T>()}, stream);
+    cudf::test::fixed_width_column_wrapper<T> 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, ApacheShortString)
 {
   auto col         = make_apache_variant(avf::short_string);
@@ -709,7 +726,10 @@ 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::FLOAT32,
+                        cudf::type_id::FLOAT64}) {
     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 bc2851ce9384ec7c8b375c6875d463c0efd218ba Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Mon, 6 Jul 2026 23:53:04 +0000
Subject: [PATCH 02/11] review changes

---
 .../parquet/experimental/variant_extract.cu   | 134 +++++-------------
 cpp/src/io/utilities/block_utils.cuh          |   4 +-
 2 files changed, 39 insertions(+), 99 deletions(-)

diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu
index 48352a723666..4acaf4e0d71d 100644
--- a/cpp/src/io/parquet/experimental/variant_extract.cu
+++ b/cpp/src/io/parquet/experimental/variant_extract.cu
@@ -376,61 +376,58 @@ constexpr bool is_variant_int =
   cuda::std::is_same_v<T, int8_t> || cuda::std::is_same_v<T, int16_t> ||
   cuda::std::is_same_v<T, int32_t> || cuda::std::is_same_v<T, int64_t>;
 
-// The floating-point types a VARIANT value can be cast to: FLOAT{32,64}.
+// The fixed-width primitive types (signed integers and floats) a VARIANT value can be decoded into.
 template <typename T>
-constexpr bool is_variant_float = cuda::std::is_same_v<T, float> || cuda::std::is_same_v<T, double>;
+constexpr bool is_variant_primitive = is_variant_int<T> || cudf::is_floating_point<T>();
 
 // The output types a VARIANT value can be cast to: the fixed-width signed integers, floats, and
 // strings.
 template <typename T>
 constexpr bool is_variant_castable =
-  is_variant_int<T> || is_variant_float<T> || cuda::std::is_same_v<T, cudf::string_view>;
+  is_variant_primitive<T> || cuda::std::is_same_v<T, cudf::string_view>;
 
-// Variant primitive ints: basic_type == primitive, value_header maps INT{8,16,32,64}.
+// Maps a fixed-width output type to the VARIANT primitive type header id that encodes it.
 template <typename T>
-__device__ inline cuda::std::optional<T> decode_int(device_span<uint8_t const> enc)
+__device__ constexpr primitive_type primitive_type_for()
 {
-  static_assert(is_variant_int<T>, "decode_int: T must be int8_t, int16_t, int32_t, or int64_t");
-
-  if (cuda::std::cmp_less(enc.size(), 1 + sizeof(T))) { return cuda::std::nullopt; }
-
-  constexpr primitive_type expected = cuda::std::is_same_v<T, int8_t>    ? primitive_type::int8
-                                      : cuda::std::is_same_v<T, int16_t> ? primitive_type::int16
-                                      : cuda::std::is_same_v<T, int32_t> ? primitive_type::int32
-                                                                         : primitive_type::int64;
-  uint8_t const value_metadata      = enc[0];
-  if (variant_basic_type(value_metadata) != basic_type::primitive ||
-      variant_value_header(value_metadata) != static_cast<uint8_t>(expected)) {
-    return cuda::std::nullopt;
+  using enum primitive_type;
+  if constexpr (cuda::std::is_same_v<T, int8_t>) {
+    return int8;
+  } else if constexpr (cuda::std::is_same_v<T, int16_t>) {
+    return int16;
+  } else if constexpr (cuda::std::is_same_v<T, int32_t>) {
+    return int32;
+  } else if constexpr (cuda::std::is_same_v<T, int64_t>) {
+    return int64;
+  } else if constexpr (cuda::std::is_same_v<T, float>) {
+    return float32;
+  } else if constexpr (cuda::std::is_same_v<T, double>) {
+    return float64;
+  } else {
+    static_assert(is_variant_primitive<T>, "primitive_type_for: no VARIANT primitive type for T");
+    return null;
   }
-  return cudf::io::unaligned_load<T>(enc.data() + 1);
 }
 
 /**
- * @brief Decode a single VARIANT value blob into a float of type `T`.
+ * @brief Decode a single VARIANT value blob into a fixed-width primitive of type `T`.
  *
- * Matches only when the blob is a primitive value whose physical type id is the exact-width float
- * header for `T` (FLOAT32 for `float`, FLOAT64 for `double`); there is no widening between float32
- * and float64. `unaligned_load` is constrained to integral types, so the payload is copied via
- * `memcpy` instead.
+ * Requires `basic_type == primitive` and a value header whose physical type id matches `T` exactly.
  */
 template <typename T>
-__device__ inline cuda::std::optional<T> decode_float(device_span<uint8_t const> enc)
+__device__ inline cuda::std::optional<T> decode_primitive(device_span<uint8_t const> enc)
 {
-  static_assert(is_variant_float<T>, "decode_float: T must be float or double");
+  static_assert(is_variant_primitive<T>,
+                "decode_primitive: T must be a variant primitive int or float");
 
   if (cuda::std::cmp_less(enc.size(), 1 + sizeof(T))) { return cuda::std::nullopt; }
 
-  constexpr primitive_type expected =
-    cuda::std::is_same_v<T, float> ? primitive_type::float32 : primitive_type::float64;
   uint8_t const value_metadata = enc[0];
   if (variant_basic_type(value_metadata) != basic_type::primitive ||
-      variant_value_header(value_metadata) != static_cast<uint8_t>(expected)) {
+      variant_value_header(value_metadata) != static_cast<uint8_t>(primitive_type_for<T>())) {
     return cuda::std::nullopt;
   }
-  T value;
-  cuda::std::memcpy(&value, enc.data() + 1, sizeof(T));
-  return value;
+  return cudf::io::unaligned_load<T>(enc.data() + 1);
 }
 
 __device__ device_span<uint8_t const> resolve_path(device_span<uint8_t const> meta,
@@ -538,15 +535,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 <typename T>
-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<T> d_output, bitmask_type* d_null_mask)
 {
   auto const num_rows = static_cast<size_type>(d_output.size());
@@ -565,45 +563,7 @@ CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_int_kernel(
     device_span<uint8_t const> const val{val_child.data<uint8_t>() + val_begin,
                                          static_cast<std::size_t>(val_end - val_begin)};
 
-    auto const decoded = decode_int<T>(val);
-    if (decoded.has_value()) {
-      d_output[row] = *decoded;
-    } else {
-      d_output[row] = 0;
-      cudf::clear_bit(d_null_mask, row);
-    }
-  }
-}
-
-/**
- * @brief Per-row kernel: decode each VARIANT value blob into a float of type `T`.
- *
- * Writes the decoded value to `d_output[row]` for non-null rows whose blob is a variant primitive
- * float whose physical type id matches `T` exactly (e.g. 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 <typename T>
-CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_float_kernel(
-  cudf::lists_column_device_view values, device_span<T> d_output, bitmask_type* d_null_mask)
-{
-  auto const num_rows = static_cast<size_type>(d_output.size());
-  auto const tid      = cudf::detail::grid_1d::global_thread_id<block_size>();
-  auto const stride   = cudf::detail::grid_1d::grid_stride<block_size>();
-
-  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);
-    auto const val_child = values.child();
-    device_span<uint8_t const> const val{val_child.data<uint8_t>() + val_begin,
-                                         static_cast<std::size_t>(val_end - val_begin)};
-
-    auto const decoded = decode_float<T>(val);
+    auto const decoded = decode_primitive<T>(val);
     if (decoded.has_value()) {
       d_output[row] = *decoded;
     } else {
@@ -677,32 +637,12 @@ struct cast_variant_fn {
 
   template <typename T>
   std::unique_ptr<column> operator()()
-    requires(is_variant_int<T>)
-  {
-    rmm::device_buffer data{num_rows * sizeof(T), stream, mr};
-
-    auto grid = cudf::detail::grid_1d{num_rows, block_size};
-    cast_variant_int_kernel<T><<<grid.num_blocks, block_size, 0, stream.value()>>>(
-      values, {static_cast<T*>(data.data()), static_cast<std::size_t>(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<column>(desired_type,
-                                    num_rows,
-                                    std::move(data),
-                                    null_count > 0 ? std::move(null_mask) : rmm::device_buffer{},
-                                    null_count);
-  }
-
-  template <typename T>
-  std::unique_ptr<column> operator()()
-    requires(is_variant_float<T>)
+    requires(is_variant_primitive<T>)
   {
     rmm::device_buffer data{num_rows * sizeof(T), stream, mr};
 
     auto grid = cudf::detail::grid_1d{num_rows, block_size};
-    cast_variant_float_kernel<T><<<grid.num_blocks, block_size, 0, stream.value()>>>(
+    cast_variant_primitive_kernel<T><<<grid.num_blocks, block_size, 0, stream.value()>>>(
       values, {static_cast<T*>(data.data()), static_cast<std::size_t>(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..958d18dacb7a 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 <typename T>
-  requires(cuda::std::is_integral_v<T>)
+  requires(cuda::std::is_integral_v<T> || cuda::std::is_floating_point_v<T>)
 inline __device__ T unaligned_load(uint8_t const* p)
 {
   T value;

From 4f09143211440be38fd5d3e40d13e7e8bb84e67c Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalekim0417@gmail.com>
Date: Fri, 10 Jul 2026 19:28:41 -0500
Subject: [PATCH 03/11] Update cpp/src/io/utilities/block_utils.cuh

Co-authored-by: Vukasin Milovanovic <vmilovanovic@nvidia.com>
---
 cpp/src/io/utilities/block_utils.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/src/io/utilities/block_utils.cuh b/cpp/src/io/utilities/block_utils.cuh
index 958d18dacb7a..d157378127d3 100644
--- a/cpp/src/io/utilities/block_utils.cuh
+++ b/cpp/src/io/utilities/block_utils.cuh
@@ -58,7 +58,7 @@ inline __device__ T warp_reduce_pos(T pos, uint32_t t)
 }
 
 template <typename T>
-  requires(cuda::std::is_integral_v<T> || cuda::std::is_floating_point_v<T>)
+  requires(cuda::std::is_trivially_copyable_v<T>)
 inline __device__ T unaligned_load(uint8_t const* p)
 {
   T value;

From fc991659235830bfcabd57895e0889467e101416 Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Mon, 13 Jul 2026 18:35:36 +0000
Subject: [PATCH 04/11] reviews

---
 cpp/src/io/parquet/experimental/variant_extract.cu | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu
index 4acaf4e0d71d..fb5fca72fe45 100644
--- a/cpp/src/io/parquet/experimental/variant_extract.cu
+++ b/cpp/src/io/parquet/experimental/variant_extract.cu
@@ -373,8 +373,7 @@ __device__ device_span<uint8_t const> locate_object_field(device_span<uint8_t co
 // exact width types (not e.g. __int128) since those are the only variant primitive int headers.
 template <typename T>
 constexpr bool is_variant_int =
-  cuda::std::is_same_v<T, int8_t> || cuda::std::is_same_v<T, int16_t> ||
-  cuda::std::is_same_v<T, int32_t> || cuda::std::is_same_v<T, int64_t>;
+  cudf::is_integral_not_bool<T>() && cudf::is_signed<T>() && !cuda::std::is_same_v<T, __int128_t>;
 
 // The fixed-width primitive types (signed integers and floats) a VARIANT value can be decoded into.
 template <typename T>
@@ -391,6 +390,8 @@ template <typename T>
 __device__ constexpr primitive_type primitive_type_for()
 {
   using enum primitive_type;
+  static_assert(is_variant_primitive<T>, "primitive_type_for: no VARIANT primitive type for T");
+
   if constexpr (cuda::std::is_same_v<T, int8_t>) {
     return int8;
   } else if constexpr (cuda::std::is_same_v<T, int16_t>) {
@@ -404,7 +405,7 @@ __device__ constexpr primitive_type primitive_type_for()
   } else if constexpr (cuda::std::is_same_v<T, double>) {
     return float64;
   } else {
-    static_assert(is_variant_primitive<T>, "primitive_type_for: no VARIANT primitive type for T");
+    CUDF_UNREACHABLE("primitive_type_for: T is not a supported variant primitive type");
     return null;
   }
 }

From 8115d8abfcabb52adfa7116112fdb88934ddf4ae Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Wed, 15 Jul 2026 19:27:20 +0000
Subject: [PATCH 05/11] cpp

---
 cpp/src/io/parquet/experimental/variant_extract.cu | 11 ++---------
 1 file changed, 2 insertions(+), 9 deletions(-)

diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu
index fb5fca72fe45..4409fad6defa 100644
--- a/cpp/src/io/parquet/experimental/variant_extract.cu
+++ b/cpp/src/io/parquet/experimental/variant_extract.cu
@@ -387,11 +387,10 @@ constexpr bool is_variant_castable =
 
 // Maps a fixed-width output type to the VARIANT primitive type header id that encodes it.
 template <typename T>
+  requires(is_variant_primitive<T>)
 __device__ constexpr primitive_type primitive_type_for()
 {
   using enum primitive_type;
-  static_assert(is_variant_primitive<T>, "primitive_type_for: no VARIANT primitive type for T");
-
   if constexpr (cuda::std::is_same_v<T, int8_t>) {
     return int8;
   } else if constexpr (cuda::std::is_same_v<T, int16_t>) {
@@ -402,11 +401,8 @@ __device__ constexpr primitive_type primitive_type_for()
     return int64;
   } else if constexpr (cuda::std::is_same_v<T, float>) {
     return float32;
-  } else if constexpr (cuda::std::is_same_v<T, double>) {
-    return float64;
   } else {
-    CUDF_UNREACHABLE("primitive_type_for: T is not a supported variant primitive type");
-    return null;
+    return float64;
   }
 }
 
@@ -418,9 +414,6 @@ __device__ constexpr primitive_type primitive_type_for()
 template <typename T>
 __device__ inline cuda::std::optional<T> decode_primitive(device_span<uint8_t const> enc)
 {
-  static_assert(is_variant_primitive<T>,
-                "decode_primitive: T must be a variant primitive int or float");
-
   if (cuda::std::cmp_less(enc.size(), 1 + sizeof(T))) { return cuda::std::nullopt; }
 
   uint8_t const value_metadata = enc[0];

From 4b65a672c39c4bf2379e490e1e8dccf62944639e Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalekim0417@gmail.com>
Date: Thu, 16 Jul 2026 16:01:58 -0500
Subject: [PATCH 06/11] Update
 cpp/src/io/parquet/experimental/variant_extract.cu

Co-authored-by: Igor Peshansky <7594381+igorpeshansky@users.noreply.github.com>
---
 cpp/src/io/parquet/experimental/variant_extract.cu | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu
index 4409fad6defa..53f852b58974 100644
--- a/cpp/src/io/parquet/experimental/variant_extract.cu
+++ b/cpp/src/io/parquet/experimental/variant_extract.cu
@@ -401,8 +401,11 @@ __device__ constexpr primitive_type primitive_type_for()
     return int64;
   } else if constexpr (cuda::std::is_same_v<T, float>) {
     return float32;
-  } else {
+  } else if constexpr (cuda::std::is_same_v<T, double>) {
     return float64;
+  } else {
+    CUDF_UNREACHABLE("primitive_type_for: T is not a supported variant primitive type");
+    return null;
   }
 }
 

From 7c7baeb3e26331ed431736aed8a43761e7c54ff0 Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Fri, 17 Jul 2026 20:01:44 +0000
Subject: [PATCH 07/11] cmake

---
 cpp/CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 0899c3cec002..7d65b74f510e 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -1461,7 +1461,7 @@ if(CUDF_BUILD_TESTUTIL)
   add_library(cudftestutil_objects OBJECT)
   target_link_libraries(
     cudftestutil_objects
-    PUBLIC cudf::cudftestutil GTest::gmock GTest::gmock_main GTest::gtest GTest::gtest_main
+    PUBLIC cudf::cudftestutil GTest::gmock GTest::gmock_main GTest::gtest GTest::gtest_main rmm::rmm
     PRIVATE cudf::cudftestutil_impl
   )
   add_library(cudf::cudftestutil_objects ALIAS cudftestutil_objects)

From 65862f67ea59501a5b520e7ca318af64da12cb4f Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Mon, 20 Jul 2026 20:47:32 +0000
Subject: [PATCH 08/11] removing change putting it in another pr

---
 cpp/CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 7d65b74f510e..0899c3cec002 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -1461,7 +1461,7 @@ if(CUDF_BUILD_TESTUTIL)
   add_library(cudftestutil_objects OBJECT)
   target_link_libraries(
     cudftestutil_objects
-    PUBLIC cudf::cudftestutil GTest::gmock GTest::gmock_main GTest::gtest GTest::gtest_main rmm::rmm
+    PUBLIC cudf::cudftestutil GTest::gmock GTest::gmock_main GTest::gtest GTest::gtest_main
     PRIVATE cudf::cudftestutil_impl
   )
   add_library(cudf::cudftestutil_objects ALIAS cudftestutil_objects)

From 5d70c9020c37189987eb5f336aef57d69e0966d7 Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Mon, 20 Jul 2026 20:52:21 +0000
Subject: [PATCH 09/11] changes

---
 .../parquet/experimental/variant_extract.cu   | 15 +++----
 .../io/experimental/variant_extract_test.cpp  | 44 ++++++++++++++-----
 2 files changed, 39 insertions(+), 20 deletions(-)

diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu
index 53f852b58974..c166ef873819 100644
--- a/cpp/src/io/parquet/experimental/variant_extract.cu
+++ b/cpp/src/io/parquet/experimental/variant_extract.cu
@@ -390,22 +390,21 @@ template <typename T>
   requires(is_variant_primitive<T>)
 __device__ constexpr primitive_type primitive_type_for()
 {
-  using enum primitive_type;
   if constexpr (cuda::std::is_same_v<T, int8_t>) {
-    return int8;
+    return primitive_type::int8;
   } else if constexpr (cuda::std::is_same_v<T, int16_t>) {
-    return int16;
+    return primitive_type::int16;
   } else if constexpr (cuda::std::is_same_v<T, int32_t>) {
-    return int32;
+    return primitive_type::int32;
   } else if constexpr (cuda::std::is_same_v<T, int64_t>) {
-    return int64;
+    return primitive_type::int64;
   } else if constexpr (cuda::std::is_same_v<T, float>) {
-    return float32;
+    return primitive_type::float32;
   } else if constexpr (cuda::std::is_same_v<T, double>) {
-    return float64;
+    return primitive_type::float64;
   } else {
     CUDF_UNREACHABLE("primitive_type_for: T is not a supported variant primitive type");
-    return null;
+    return primitive_type::null;
   }
 }
 
diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp
index 157605272878..9fecd299c804 100644
--- a/cpp/tests/io/experimental/variant_extract_test.cpp
+++ b/cpp/tests/io/experimental/variant_extract_test.cpp
@@ -646,16 +646,30 @@ TEST_F(CastVariantTest, ApachePrimitiveInts)
     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(
+    return cudf::io::parquet::experimental::cast_variant(
       value, cudf::data_type{cudf::type_to_id<T>()}, stream);
-    cudf::test::fixed_width_column_wrapper<T> expected{expected_val};
-    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   };
 
-  cast(avf::primitive_int8, int8_t{42});
-  cast(avf::primitive_int16, int16_t{1234});
-  cast(avf::primitive_int32, int32_t{123456});
-  cast(avf::primitive_int64, int64_t{1234567890123456789LL});
+  {
+    auto got = cast(avf::primitive_int8, int8_t{42});
+    cudf::test::fixed_width_column_wrapper<int8_t> expected{int8_t{42}};
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
+  }
+  {
+    auto got = cast(avf::primitive_int16, int16_t{1234});
+    cudf::test::fixed_width_column_wrapper<int16_t> expected{int16_t{1234}};
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
+  }
+  {
+    auto got = cast(avf::primitive_int32, int32_t{123456});
+    cudf::test::fixed_width_column_wrapper<int32_t> expected{int32_t{123456}};
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
+  }
+  {
+    auto got = cast(avf::primitive_int64, int64_t{1234567890123456789LL});
+    cudf::test::fixed_width_column_wrapper<int64_t> expected{int64_t{1234567890123456789LL}};
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
+  }
 }
 
 TEST_F(CastVariantTest, ApachePrimitiveFloats)
@@ -665,14 +679,20 @@ TEST_F(CastVariantTest, ApachePrimitiveFloats)
     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(
+    return cudf::io::parquet::experimental::cast_variant(
       value, cudf::data_type{cudf::type_to_id<T>()}, stream);
-    cudf::test::fixed_width_column_wrapper<T> expected{expected_val};
-    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   };
 
-  cast(avf::primitive_float, float{1234567936.0f});
-  cast(avf::primitive_double, double{1234567890.1234});
+  {
+    auto got = cast(avf::primitive_float, float{1234567936.0f});
+    cudf::test::fixed_width_column_wrapper<float> expected{float{1234567936.0f}};
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
+  }
+  {
+    auto got = cast(avf::primitive_double, double{1234567890.1234});
+    cudf::test::fixed_width_column_wrapper<double> expected{double{1234567890.1234}};
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
+  }
 }
 
 TEST_F(CastVariantTest, ApacheShortString)

From 9030bda9795e47e47b34ba7ddbac104f898d8308 Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Tue, 21 Jul 2026 04:04:09 +0000
Subject: [PATCH 10/11] comments

---
 .../io/experimental/variant_extract_test.cpp  | 40 ++++++++-----------
 1 file changed, 17 insertions(+), 23 deletions(-)

diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp
index 9fecd299c804..e15893c41a47 100644
--- a/cpp/tests/io/experimental/variant_extract_test.cpp
+++ b/cpp/tests/io/experimental/variant_extract_test.cpp
@@ -637,36 +637,38 @@ TEST_F(GetVariantFieldTest, EmptyInput)
   EXPECT_EQ(cudf::lists_column_view{got->view()}.child().type().id(), cudf::type_id::UINT8);
 }
 
-struct CastVariantTest : public cudf::test::BaseFixture {};
-
-TEST_F(CastVariantTest, ApachePrimitiveInts)
-{
-  auto stream     = cudf::test::get_default_stream();
-  auto const cast = [&](auto const& fixture, auto expected_val) {
-    using T          = decltype(expected_val);
+struct CastVariantTest : public cudf::test::BaseFixture {
+  template <typename T, std::size_t M, std::size_t V>
+  std::unique_ptr<cudf::column> cast_apache_primitive(avf::fixture<M, V> const& fixture,
+                                                      rmm::cuda_stream_view 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::cast_variant(
       value, cudf::data_type{cudf::type_to_id<T>()}, stream);
-  };
+  }
+};
 
+TEST_F(CastVariantTest, ApachePrimitiveInts)
+{
+  auto stream = cudf::test::get_default_stream();
   {
-    auto got = cast(avf::primitive_int8, int8_t{42});
+    auto got = cast_apache_primitive<int8_t>(avf::primitive_int8, stream);
     cudf::test::fixed_width_column_wrapper<int8_t> expected{int8_t{42}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast(avf::primitive_int16, int16_t{1234});
+    auto got = cast_apache_primitive<int16_t>(avf::primitive_int16, stream);
     cudf::test::fixed_width_column_wrapper<int16_t> expected{int16_t{1234}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast(avf::primitive_int32, int32_t{123456});
+    auto got = cast_apache_primitive<int32_t>(avf::primitive_int32, stream);
     cudf::test::fixed_width_column_wrapper<int32_t> expected{int32_t{123456}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast(avf::primitive_int64, int64_t{1234567890123456789LL});
+    auto got = cast_apache_primitive<int64_t>(avf::primitive_int64, stream);
     cudf::test::fixed_width_column_wrapper<int64_t> expected{int64_t{1234567890123456789LL}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
@@ -674,22 +676,14 @@ TEST_F(CastVariantTest, ApachePrimitiveInts)
 
 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);
-    return cudf::io::parquet::experimental::cast_variant(
-      value, cudf::data_type{cudf::type_to_id<T>()}, stream);
-  };
-
+  auto stream = cudf::test::get_default_stream();
   {
-    auto got = cast(avf::primitive_float, float{1234567936.0f});
+    auto got = cast_apache_primitive<float>(avf::primitive_float, stream);
     cudf::test::fixed_width_column_wrapper<float> expected{float{1234567936.0f}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast(avf::primitive_double, double{1234567890.1234});
+    auto got = cast_apache_primitive<double>(avf::primitive_double, stream);
     cudf::test::fixed_width_column_wrapper<double> expected{double{1234567890.1234}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }

From f2efc367e9798b52e0f655a31cd9ccf1512f8475 Mon Sep 17 00:00:00 2001
From: Abigale Kim <abigalek@nvidia.com>
Date: Wed, 22 Jul 2026 17:30:03 +0000
Subject: [PATCH 11/11] cpp

---
 .../io/experimental/variant_extract_test.cpp  | 36 +++++++++----------
 1 file changed, 17 insertions(+), 19 deletions(-)

diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp
index e15893c41a47..2e53fd92bceb 100644
--- a/cpp/tests/io/experimental/variant_extract_test.cpp
+++ b/cpp/tests/io/experimental/variant_extract_test.cpp
@@ -637,38 +637,37 @@ TEST_F(GetVariantFieldTest, EmptyInput)
   EXPECT_EQ(cudf::lists_column_view{got->view()}.child().type().id(), cudf::type_id::UINT8);
 }
 
-struct CastVariantTest : public cudf::test::BaseFixture {
-  template <typename T, std::size_t M, std::size_t V>
-  std::unique_ptr<cudf::column> cast_apache_primitive(avf::fixture<M, V> const& fixture,
-                                                      rmm::cuda_stream_view 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::cast_variant(
-      value, cudf::data_type{cudf::type_to_id<T>()}, stream);
-  }
-};
+template <typename T, std::size_t M, std::size_t V>
+std::unique_ptr<cudf::column> cast_apache_primitive(avf::fixture<M, V> 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::cast_variant(
+    value, cudf::data_type{cudf::type_to_id<T>()}, stream);
+}
+
+struct CastVariantTest : public cudf::test::BaseFixture {};
 
 TEST_F(CastVariantTest, ApachePrimitiveInts)
 {
-  auto stream = cudf::test::get_default_stream();
   {
-    auto got = cast_apache_primitive<int8_t>(avf::primitive_int8, stream);
+    auto got = cast_apache_primitive<int8_t>(avf::primitive_int8);
     cudf::test::fixed_width_column_wrapper<int8_t> expected{int8_t{42}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast_apache_primitive<int16_t>(avf::primitive_int16, stream);
+    auto got = cast_apache_primitive<int16_t>(avf::primitive_int16);
     cudf::test::fixed_width_column_wrapper<int16_t> expected{int16_t{1234}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast_apache_primitive<int32_t>(avf::primitive_int32, stream);
+    auto got = cast_apache_primitive<int32_t>(avf::primitive_int32);
     cudf::test::fixed_width_column_wrapper<int32_t> expected{int32_t{123456}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast_apache_primitive<int64_t>(avf::primitive_int64, stream);
+    auto got = cast_apache_primitive<int64_t>(avf::primitive_int64);
     cudf::test::fixed_width_column_wrapper<int64_t> expected{int64_t{1234567890123456789LL}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
@@ -676,14 +675,13 @@ TEST_F(CastVariantTest, ApachePrimitiveInts)
 
 TEST_F(CastVariantTest, ApachePrimitiveFloats)
 {
-  auto stream = cudf::test::get_default_stream();
   {
-    auto got = cast_apache_primitive<float>(avf::primitive_float, stream);
+    auto got = cast_apache_primitive<float>(avf::primitive_float);
     cudf::test::fixed_width_column_wrapper<float> expected{float{1234567936.0f}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }
   {
-    auto got = cast_apache_primitive<double>(avf::primitive_double, stream);
+    auto got = cast_apache_primitive<double>(avf::primitive_double);
     cudf::test::fixed_width_column_wrapper<double> expected{double{1234567890.1234}};
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected);
   }