From 5b3e7e783a625b68cf0e2675a637827ced875ba2 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 25 Feb 2026 10:36:02 +0000
Subject: [PATCH 01/19] [FEA] Support dictionary-encoded types in Transforms

---
 .../cudf/column/column_device_view_base.cuh   | 45 +++++++++-
 cpp/src/jit/helpers.cpp                       | 30 +++++--
 cpp/src/transform/transform.cu                | 14 ++--
 .../integration/unary_transform_test.cpp      | 84 +++++++++++++++++++
 4 files changed, 161 insertions(+), 12 deletions(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 111361b10961..2269c4574eae 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -63,6 +63,19 @@ struct nullate {
   };
 };
 
+template <typename IndexType, typename KeyType>
+  requires(is_index_type<IndexType>())
+struct dictionary_element {
+  using index_type = IndexType;
+  using key_type   = KeyType;
+};
+
+template <typename T>
+inline constexpr bool is_dictionary_encoded = false;
+
+template <typename IndexType, typename KeyType>
+inline constexpr bool is_dictionary_encoded<dictionary_element<IndexType, KeyType>> = true;
+
 namespace detail {
 /**
  * @brief An immutable, non-owning view of device data as a column of elements
@@ -79,6 +92,10 @@ class alignas(16) column_device_view_base {
  public:
   // TODO: merge this offsets column index with `strings_column_view::offsets_column_index`
   static constexpr size_type offsets_column_index{0};  ///< Child index of the offsets column
+  static constexpr size_type dictionary_offsets_column_index =
+    0;  ///< Child index of the dictionary offsets column
+  static constexpr size_type dictionary_keys_column_index =
+    1;  ///< Child index of the dictionary key column
 
   column_device_view_base()                               = delete;
   ~column_device_view_base()                              = default;
@@ -453,6 +470,32 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
     return T{scaled_integer<rep>{data<rep>()[element_index], scale}};
   }
 
+  /**
+   * @brief Returns a copy of the element at the specified index
+   *
+   * If the element at the specified index is NULL, i.e.,
+   * `is_null(element_index) == true`, then any attempt to use the result will
+   * lead to undefined behavior.
+   *
+   * This function accounts for the offset.
+   *
+   * This function does not participate in overload resolution if `is_dictionary_encoded<T>` is
+   * false.
+   *
+   * @tparam T The element type, i.e. dictionary_element<int32_t, float> for a dictionary column
+   * with int32 indices and float values
+   * @param element_index Position of the desired element
+   * @return The element at the specified index
+   */
+  template <typename T, CUDF_ENABLE_IF(is_dictionary_encoded<T>)>
+  [[nodiscard]] __device__ decltype(auto) element(size_type element_index) const noexcept
+  {
+    auto const& offsets = child(dictionary_offsets_column_index);
+    auto const& keys    = child(dictionary_keys_column_index);
+    auto const index    = offsets.template element<typename T::index_type>(element_index);
+    return keys.template element<typename T::key_type>(index);
+  }
+
   /**
    * @brief Returns the specified child
    *
diff --git a/cpp/src/jit/helpers.cpp b/cpp/src/jit/helpers.cpp
index 9541950ad785..9626634e0847 100644
--- a/cpp/src/jit/helpers.cpp
+++ b/cpp/src/jit/helpers.cpp
@@ -5,10 +5,13 @@
 
 #include "helpers.hpp"
 
+#include <cudf/column/column_device_view_base.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 
 #include <jit/cache.hpp>
 
+#include <format>
+
 namespace cudf {
 namespace jit {
 
@@ -111,13 +114,30 @@ std::vector<std::string> input_type_names(
   return names;
 }
 
-input_reflection reflect_input(std::variant<column_view, scalar_column_view> const& input)
+std::string get_jit_element_type_name(column_view const& view)
 {
-  auto get_type_name = [](auto const& var) {
-    return std::visit([](auto& a) { return type_to_name(a.type()); }, var);
-  };
+  if (is_fixed_width(view.type()) || view.type().id() == type_id::STRING) {
+    return type_to_name(view.type());
+  } else if (view.type().id() == type_id::DICTIONARY32) {
+    return std::format(
+      "cudf::dictionary_element<{}, {}>",
+      get_jit_element_type_name(
+        view.child(column_device_view_core::dictionary_offsets_column_index)),
+      get_jit_element_type_name(view.child(column_device_view_core::dictionary_keys_column_index)));
+  } else {
+    CUDF_FAIL("Unsupported type for JIT compilation: " + type_to_name(view.type()));
+  }
+}
+
+std::string get_jit_element_type_name(scalar_column_view const& view)
+{
+  return get_jit_element_type_name(view.as_column_view());
+}
 
-  return input_reflection{get_type_name(input), std::holds_alternative<scalar_column_view>(input)};
+input_reflection reflect_input(std::variant<column_view, scalar_column_view> const& input)
+{
+  auto type_name = std::visit([](auto& a) { return get_jit_element_type_name(a); }, input);
+  return input_reflection{type_name, std::holds_alternative<scalar_column_view>(input)};
 }
 
 std::vector<input_reflection> reflect_inputs(
diff --git a/cpp/src/transform/transform.cu b/cpp/src/transform/transform.cu
index 9dd362f6e26c..61399753c9bf 100644
--- a/cpp/src/transform/transform.cu
+++ b/cpp/src/transform/transform.cu
@@ -425,12 +425,14 @@ void perform_checks(std::optional<size_type> in_row_size, data_type output_type,
     return std::visit([](auto const& col) { return col.type(); }, in);
   };
 
-  CUDF_EXPECTS(
-    std::all_of(thrust::make_transform_iterator(inputs.begin(), get_type),
-                thrust::make_transform_iterator(inputs.end(), get_type),
-                [](data_type t) { return is_fixed_width(t) || (t.id() == type_id::STRING); }),
-    "Transforms only support input of fixed-width or string types",
-    std::invalid_argument);
+  CUDF_EXPECTS(std::all_of(thrust::make_transform_iterator(inputs.begin(), get_type),
+                           thrust::make_transform_iterator(inputs.end(), get_type),
+                           [](data_type t) {
+                             return is_fixed_width(t) || (t.id() == type_id::STRING) ||
+                                    (t.id() == type_id::DICTIONARY32);
+                           }),
+               "Transforms only support input of fixed-width, string, and dictionary types",
+               std::invalid_argument);
 
   check_row_size(in_row_size, inputs);
 }
diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp
index 49720a27c848..19584ed1ba97 100644
--- a/cpp/tests/transform/integration/unary_transform_test.cpp
+++ b/cpp/tests/transform/integration/unary_transform_test.cpp
@@ -29,6 +29,7 @@
 #include <cudf_test/type_lists.hpp>
 
 #include <cudf/detail/iterator.cuh>
+#include <cudf/dictionary/encode.hpp>
 #include <cudf/transform.hpp>
 
 namespace transformation {
@@ -309,6 +310,89 @@ __device__ inline void f(cudf::timestamp_us* output, cudf::timestamp_us input)
   test_udf<dtype>(cuda.c_str(), op, data_init, 500, false);
 }
 
+TEST_F(UnaryOperationIntegrationTest, Transform_DictionaryString)
+{
+  std::string const cuda =
+    R"***(
+__device__ inline void decode(cudf::string_view * output, cudf::string_view input){
+  *output = input;
+})***";
+
+  // non-nullable
+  {
+    auto a =
+      cudf::test::strings_column_wrapper{
+        "eee", "aaa", "ddd", "bbb", "ccc", "ccc", "ccc", "eee", "aaa"}
+        .release();
+
+    auto a_encoded = cudf::dictionary::encode(a->view());
+
+    cudf::transform_input inputs[] = {*a_encoded};
+
+    auto out =
+      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::STRING}, false);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
+  }
+
+  // nullable
+  {
+    auto a =
+      cudf::test::strings_column_wrapper{
+        {"eee", "aaa", "ddd", "bbb", "ccc", "ccc", "ccc", "eee", "aaa"},
+        {true, true, true, false, true, true, true, true, true}}
+        .release();
+
+    auto a_encoded = cudf::dictionary::encode(a->view());
+
+    cudf::transform_input inputs[] = {*a_encoded};
+
+    auto out =
+      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::STRING}, false);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
+  }
+}
+
+TEST_F(UnaryOperationIntegrationTest, Transform_DictionaryFloat)
+{
+  std::string const cuda =
+    R"***(
+__device__ inline void decode(float * output, float input){
+  *output = input;
+})***";
+
+  // non-nullable
+  {
+    auto a = cudf::test::fixed_width_column_wrapper<float>(
+               {1.0F, 2.0F, 3.0F, 4.0F, 5.0F, 5.0F, 5.0F, 1.0F, 2.0F})
+               .release();
+    auto a_encoded                 = cudf::dictionary::encode(a->view());
+    cudf::transform_input inputs[] = {*a_encoded};
+
+    auto out =
+      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, false);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
+  }
+
+  // nullable
+  {
+    auto a = cudf::test::fixed_width_column_wrapper<float>(
+               {{1.0F, 2.0F, 3.0F, 4.0F, 5.0F, 5.0F, 5.0F, 1.0F, 2.0F},
+                {true, true, true, true, true, false, true, true, true}})
+               .release();
+
+    auto a_encoded                 = cudf::dictionary::encode(a->view());
+    cudf::transform_input inputs[] = {*a_encoded};
+
+    auto out =
+      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, false);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
+  }
+}
+
 struct TernaryOperationTest : public cudf::test::BaseFixture {};
 
 TEST_F(TernaryOperationTest, TransformWithScalar)

From ebb44190f19eaeb30159089f63a3e90eacb4e480 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 25 Feb 2026 10:49:38 +0000
Subject: [PATCH 02/19] [FEA] Enhance dictionary_element with key member and
 improve documentation

---
 .../cudf/column/column_device_view_base.cuh   | 20 ++++++++++++++-----
 1 file changed, 15 insertions(+), 5 deletions(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 2269c4574eae..61be0e650627 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -63,13 +63,23 @@ struct nullate {
   };
 };
 
+/**
+ * @brief A type tag to specify that a column should be treated as a dictionary column.
+ * @tparam IndexType The type of the dictionary indices
+ * @tparam KeyType The type of the dictionary keys
+ */
 template <typename IndexType, typename KeyType>
   requires(is_index_type<IndexType>())
 struct dictionary_element {
-  using index_type = IndexType;
-  using key_type   = KeyType;
+  using index_type = IndexType;  ///< The type of the dictionary indices
+  using key_type   = KeyType;    ///< The type of the dictionary keys
+
+  key_type key{};  ///< The dictionary key for this element
 };
 
+/**
+ * @brief A type trait to determine if a type is a dictionary encoded type.
+ */
 template <typename T>
 inline constexpr bool is_dictionary_encoded = false;
 
@@ -471,7 +481,7 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
   }
 
   /**
-   * @brief Returns a copy of the element at the specified index
+   * @brief Returns a decoded copy of the element at the specified index.
    *
    * If the element at the specified index is NULL, i.e.,
    * `is_null(element_index) == true`, then any attempt to use the result will
@@ -482,8 +492,8 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
    * This function does not participate in overload resolution if `is_dictionary_encoded<T>` is
    * false.
    *
-   * @tparam T The element type, i.e. dictionary_element<int32_t, float> for a dictionary column
-   * with int32 indices and float values
+   * @tparam T The element type, i.e. `dictionary_element<int32_t, float>` for a dictionary column
+   * with `int32_t` indices and `float` values
    * @param element_index Position of the desired element
    * @return The element at the specified index
    */

From eb6b24d3a5360515f74ca95334812acfa743a3d4 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 25 Feb 2026 10:59:24 +0000
Subject: [PATCH 03/19] fix style check

---
 cpp/include/cudf/column/column_device_view_base.cuh | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 61be0e650627..84b3cc9191c1 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -79,10 +79,16 @@ struct dictionary_element {
 
 /**
  * @brief A type trait to determine if a type is a dictionary encoded type.
+ * @tparam T The type to check
  */
 template <typename T>
 inline constexpr bool is_dictionary_encoded = false;
 
+/**
+ * @brief A type trait to determine if a type is a dictionary encoded type.
+ * @tparam IndexType The type of the dictionary indices
+ * @tparam KeyType The type of the dictionary keys
+ */
 template <typename IndexType, typename KeyType>
 inline constexpr bool is_dictionary_encoded<dictionary_element<IndexType, KeyType>> = true;
 

From 74db1a8fb48797744b5aa59638aef6583c11b10a Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 25 Feb 2026 14:09:55 +0000
Subject: [PATCH 04/19] code review changes

---
 cpp/src/jit/helpers.cpp        | 46 ++++++++++++++++++++++++++++------
 cpp/src/transform/transform.cu |  2 +-
 2 files changed, 39 insertions(+), 9 deletions(-)

diff --git a/cpp/src/jit/helpers.cpp b/cpp/src/jit/helpers.cpp
index 9626634e0847..ce8c74deb0b6 100644
--- a/cpp/src/jit/helpers.cpp
+++ b/cpp/src/jit/helpers.cpp
@@ -7,6 +7,7 @@
 
 #include <cudf/column/column_device_view_base.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/utilities/type_dispatcher.hpp>
 
 #include <jit/cache.hpp>
 
@@ -114,19 +115,46 @@ std::vector<std::string> input_type_names(
   return names;
 }
 
-std::string get_jit_element_type_name(column_view const& view)
-{
-  if (is_fixed_width(view.type()) || view.type().id() == type_id::STRING) {
+namespace {
+
+std::string get_jit_element_type_name_impl(column_view const& view);
+
+struct jit_element_type_name_fn {
+  template <typename T>
+    requires(is_fixed_width<T>() || std::is_same_v<T, cudf::string_view>)
+  std::string operator()(column_view const& view) const
+  {
     return type_to_name(view.type());
-  } else if (view.type().id() == type_id::DICTIONARY32) {
+  }
+
+  template <typename T>
+    requires(std::is_same_v<T, cudf::dictionary32>)
+  std::string operator()(column_view const& view) const
+  {
     return std::format(
       "cudf::dictionary_element<{}, {}>",
-      get_jit_element_type_name(
-        view.child(column_device_view_core::dictionary_offsets_column_index)),
-      get_jit_element_type_name(view.child(column_device_view_core::dictionary_keys_column_index)));
-  } else {
+      get_jit_element_type_name_impl(
+        view.child(column_device_view::dictionary_indices_column_index)),
+      get_jit_element_type_name_impl(view.child(column_device_view::dictionary_keys_column_index)));
+  }
+
+  template <typename T>
+    requires(!is_fixed_width<T>() && !std::is_same_v<T, cudf::string_view> &&
+             !std::is_same_v<T, cudf::dictionary32>)
+  std::string operator()(column_view const& view) const
+  {
     CUDF_FAIL("Unsupported type for JIT compilation: " + type_to_name(view.type()));
   }
+};
+
+std::string get_jit_element_type_name_impl(column_view const& view)
+{
+  return cudf::type_dispatcher(view.type(), jit_element_type_name_fn{}, view);
+}
+
+std::string get_jit_element_type_name(column_view const& view)
+{
+  return get_jit_element_type_name_impl(view);
 }
 
 std::string get_jit_element_type_name(scalar_column_view const& view)
@@ -134,6 +162,8 @@ std::string get_jit_element_type_name(scalar_column_view const& view)
   return get_jit_element_type_name(view.as_column_view());
 }
 
+}  // namespace
+
 input_reflection reflect_input(std::variant<column_view, scalar_column_view> const& input)
 {
   auto type_name = std::visit([](auto& a) { return get_jit_element_type_name(a); }, input);
diff --git a/cpp/src/transform/transform.cu b/cpp/src/transform/transform.cu
index 61399753c9bf..f41b2cae1c97 100644
--- a/cpp/src/transform/transform.cu
+++ b/cpp/src/transform/transform.cu
@@ -429,7 +429,7 @@ void perform_checks(std::optional<size_type> in_row_size, data_type output_type,
                            thrust::make_transform_iterator(inputs.end(), get_type),
                            [](data_type t) {
                              return is_fixed_width(t) || (t.id() == type_id::STRING) ||
-                                    (t.id() == type_id::DICTIONARY32);
+                                    cudf::is_dictionary(t);
                            }),
                "Transforms only support input of fixed-width, string, and dictionary types",
                std::invalid_argument);

From c298ff55230643bc39d4fca1565dd7f47b42b16e Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 25 Feb 2026 21:39:22 +0000
Subject: [PATCH 05/19] code review changes

---
 cpp/include/cudf/column/column_device_view_base.cuh    | 8 ++++----
 cpp/include/cudf/dictionary/dictionary_column_view.hpp | 6 ++++--
 2 files changed, 8 insertions(+), 6 deletions(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 84b3cc9191c1..71c8c5237aed 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -65,6 +65,7 @@ struct nullate {
 
 /**
  * @brief A type tag to specify that a column should be treated as a dictionary column.
+ *
  * @tparam IndexType The type of the dictionary indices
  * @tparam KeyType The type of the dictionary keys
  */
@@ -106,9 +107,8 @@ namespace detail {
  */
 class alignas(16) column_device_view_base {
  public:
-  // TODO: merge this offsets column index with `strings_column_view::offsets_column_index`
   static constexpr size_type offsets_column_index{0};  ///< Child index of the offsets column
-  static constexpr size_type dictionary_offsets_column_index =
+  static constexpr size_type dictionary_indices_column_index =
     0;  ///< Child index of the dictionary offsets column
   static constexpr size_type dictionary_keys_column_index =
     1;  ///< Child index of the dictionary key column
@@ -506,9 +506,9 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
   template <typename T, CUDF_ENABLE_IF(is_dictionary_encoded<T>)>
   [[nodiscard]] __device__ decltype(auto) element(size_type element_index) const noexcept
   {
-    auto const& offsets = child(dictionary_offsets_column_index);
+    auto const& indices = child(dictionary_indices_column_index);
     auto const& keys    = child(dictionary_keys_column_index);
-    auto const index    = offsets.template element<typename T::index_type>(element_index);
+    auto const index    = indices.template element<typename T::index_type>(element_index);
     return keys.template element<typename T::key_type>(index);
   }
 
diff --git a/cpp/include/cudf/dictionary/dictionary_column_view.hpp b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
index 2c4659a508c8..b82b08d24ae8 100644
--- a/cpp/include/cudf/dictionary/dictionary_column_view.hpp
+++ b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
@@ -4,6 +4,7 @@
  */
 #pragma once
 
+#include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_view.hpp>
 
 /**
@@ -52,9 +53,10 @@ class dictionary_column_view : private column_view {
   dictionary_column_view& operator=(dictionary_column_view&&) = default;
 
   /// Index of the indices column of the dictionary column
-  static constexpr size_type indices_column_index{0};
+  static constexpr size_type indices_column_index =
+    column_device_view::dictionary_indices_column_index;
   /// Index of the keys column of the dictionary column
-  static constexpr size_type keys_column_index{1};
+  static constexpr size_type keys_column_index = column_device_view::dictionary_keys_column_index;
 
   using column_view::has_nulls;
   using column_view::is_empty;

From 70cc26ba80993714b895e5e22367a5fd95c4df15 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 25 Feb 2026 21:45:16 +0000
Subject: [PATCH 06/19] Update
 cpp/include/cudf/column/column_device_view_base.cuh

Co-authored-by: David Wendt <45795991+davidwendt@users.noreply.github.com>
---
 cpp/include/cudf/column/column_device_view_base.cuh | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 71c8c5237aed..b69c37b37175 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -65,7 +65,8 @@ struct nullate {
 
 /**
  * @brief A type tag to specify that a column should be treated as a dictionary column.
- *
+ * @brief A type tag to specify that a column should be treated as a dictionary column
+ * 
  * @tparam IndexType The type of the dictionary indices
  * @tparam KeyType The type of the dictionary keys
  */

From 15e8cc02e21a12fe03a56859e25ee0878b16fc12 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Fri, 27 Feb 2026 13:40:10 +0000
Subject: [PATCH 07/19] update

---
 .../cudf/column/column_child_offsets.h        | 20 +++++++++++++++++++
 .../cudf/column/column_device_view_base.cuh   | 12 +++++------
 .../dictionary/dictionary_column_view.hpp     |  9 ++++-----
 cpp/src/jit/helpers.cpp                       |  5 ++---
 4 files changed, 31 insertions(+), 15 deletions(-)
 create mode 100644 cpp/include/cudf/column/column_child_offsets.h

diff --git a/cpp/include/cudf/column/column_child_offsets.h b/cpp/include/cudf/column/column_child_offsets.h
new file mode 100644
index 000000000000..4a96f87c6835
--- /dev/null
+++ b/cpp/include/cudf/column/column_child_offsets.h
@@ -0,0 +1,20 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/types.hpp>
+#include <cudf/utilities/export.hpp>
+
+namespace CUDF_EXPORT cudf {
+
+static constexpr size_type offsets_column_index = 0;  ///< Child index of the offsets column
+
+static constexpr size_type dictionary_indices_column_index =
+  0;  ///< Child index of the dictionary offsets column
+
+static constexpr size_type dictionary_keys_column_index =
+  1;  ///< Child index of the dictionary key column
+
+}  // namespace CUDF_EXPORT cudf
\ No newline at end of file
diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index b69c37b37175..de0fad6914bc 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -4,6 +4,7 @@
  */
 #pragma once
 
+#include <cudf/column/column_child_offsets.h>
 #include <cudf/detail/offsets_iterator.cuh>
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/strings/string_view.cuh>
@@ -66,12 +67,12 @@ struct nullate {
 /**
  * @brief A type tag to specify that a column should be treated as a dictionary column.
  * @brief A type tag to specify that a column should be treated as a dictionary column
- * 
+ *
  * @tparam IndexType The type of the dictionary indices
  * @tparam KeyType The type of the dictionary keys
  */
 template <typename IndexType, typename KeyType>
-  requires(is_index_type<IndexType>())
+  requires(is_index_type<IndexType>() && is_relationally_comparable<KeyType, KeyType>())
 struct dictionary_element {
   using index_type = IndexType;  ///< The type of the dictionary indices
   using key_type   = KeyType;    ///< The type of the dictionary keys
@@ -108,11 +109,8 @@ namespace detail {
  */
 class alignas(16) column_device_view_base {
  public:
-  static constexpr size_type offsets_column_index{0};  ///< Child index of the offsets column
-  static constexpr size_type dictionary_indices_column_index =
-    0;  ///< Child index of the dictionary offsets column
-  static constexpr size_type dictionary_keys_column_index =
-    1;  ///< Child index of the dictionary key column
+  static constexpr size_type offsets_column_index =
+    cudf::offsets_column_index;  ///< Child index of the offsets column
 
   column_device_view_base()                               = delete;
   ~column_device_view_base()                              = default;
diff --git a/cpp/include/cudf/dictionary/dictionary_column_view.hpp b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
index b82b08d24ae8..8666e0ec008d 100644
--- a/cpp/include/cudf/dictionary/dictionary_column_view.hpp
+++ b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
@@ -1,10 +1,10 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
 
-#include <cudf/column/column_device_view.cuh>
+#include <cudf/column/column_child_offsets.h>
 #include <cudf/column/column_view.hpp>
 
 /**
@@ -53,10 +53,9 @@ class dictionary_column_view : private column_view {
   dictionary_column_view& operator=(dictionary_column_view&&) = default;
 
   /// Index of the indices column of the dictionary column
-  static constexpr size_type indices_column_index =
-    column_device_view::dictionary_indices_column_index;
+  static constexpr size_type indices_column_index = cudf::dictionary_indices_column_index;
   /// Index of the keys column of the dictionary column
-  static constexpr size_type keys_column_index = column_device_view::dictionary_keys_column_index;
+  static constexpr size_type keys_column_index = cudf::dictionary_keys_column_index;
 
   using column_view::has_nulls;
   using column_view::is_empty;
diff --git a/cpp/src/jit/helpers.cpp b/cpp/src/jit/helpers.cpp
index ce8c74deb0b6..b8b5c7d2a941 100644
--- a/cpp/src/jit/helpers.cpp
+++ b/cpp/src/jit/helpers.cpp
@@ -133,9 +133,8 @@ struct jit_element_type_name_fn {
   {
     return std::format(
       "cudf::dictionary_element<{}, {}>",
-      get_jit_element_type_name_impl(
-        view.child(column_device_view::dictionary_indices_column_index)),
-      get_jit_element_type_name_impl(view.child(column_device_view::dictionary_keys_column_index)));
+      get_jit_element_type_name_impl(view.child(cudf::dictionary_indices_column_index)),
+      get_jit_element_type_name_impl(view.child(cudf::dictionary_keys_column_index)));
   }
 
   template <typename T>

From 3553432e35e9969211da154cb49369e68b781b38 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Fri, 27 Feb 2026 13:40:33 +0000
Subject: [PATCH 08/19] update

---
 cpp/include/cudf/column/column_child_offsets.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/include/cudf/column/column_child_offsets.h b/cpp/include/cudf/column/column_child_offsets.h
index 4a96f87c6835..27ebe11e3e6c 100644
--- a/cpp/include/cudf/column/column_child_offsets.h
+++ b/cpp/include/cudf/column/column_child_offsets.h
@@ -17,4 +17,4 @@ static constexpr size_type dictionary_indices_column_index =
 static constexpr size_type dictionary_keys_column_index =
   1;  ///< Child index of the dictionary key column
 
-}  // namespace CUDF_EXPORT cudf
\ No newline at end of file
+}  // namespace CUDF_EXPORT cudf

From 10d2ebd707c32ce075df504269ad2fcc2ff68a7c Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Fri, 27 Feb 2026 17:37:05 +0000
Subject: [PATCH 09/19] Update
 cpp/include/cudf/column/column_device_view_base.cuh

Co-authored-by: Bradley Dice <bdice@bradleydice.com>
---
 cpp/include/cudf/column/column_device_view_base.cuh | 1 -
 1 file changed, 1 deletion(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index de0fad6914bc..edce19fbaa06 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -65,7 +65,6 @@ struct nullate {
 };
 
 /**
- * @brief A type tag to specify that a column should be treated as a dictionary column.
  * @brief A type tag to specify that a column should be treated as a dictionary column
  *
  * @tparam IndexType The type of the dictionary indices

From 9cea60a930df7b1f941d75271b3735150262c2ff Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Wed, 4 Mar 2026 22:02:58 +0000
Subject: [PATCH 10/19] Update
 cpp/include/cudf/column/column_device_view_base.cuh

Co-authored-by: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
---
 cpp/include/cudf/column/column_device_view_base.cuh | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 3036e6501ea2..0dea689de5ea 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -490,8 +490,8 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
    * This function does not participate in overload resolution if `is_dictionary_encoded<T>` is
    * false.
    *
-   * @tparam T The element type, i.e. `dictionary_element<int32_t, float>` for a dictionary column
-   * with `int32_t` indices and `float` values
+   * @tparam T The element type, e.g., `dictionary_element<int32_t, float>` for a dictionary column
+   * with `int32_t` indices and `float` keys
    * @param element_index Position of the desired element
    * @return The element at the specified index
    */

From de00c9a33793c3ef8711fb704bd7dfecc1b11184 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Fri, 13 Mar 2026 08:26:04 +0000
Subject: [PATCH 11/19] code review changes

---
 .../{column_child_offsets.h => column_child_offsets.hpp}      | 2 +-
 cpp/include/cudf/transform.hpp                                | 4 ++--
 cpp/tests/transform/integration/unary_transform_test.cpp      | 4 ++--
 3 files changed, 5 insertions(+), 5 deletions(-)
 rename cpp/include/cudf/column/{column_child_offsets.h => column_child_offsets.hpp} (87%)

diff --git a/cpp/include/cudf/column/column_child_offsets.h b/cpp/include/cudf/column/column_child_offsets.hpp
similarity index 87%
rename from cpp/include/cudf/column/column_child_offsets.h
rename to cpp/include/cudf/column/column_child_offsets.hpp
index 27ebe11e3e6c..0b9435f94338 100644
--- a/cpp/include/cudf/column/column_child_offsets.h
+++ b/cpp/include/cudf/column/column_child_offsets.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
diff --git a/cpp/include/cudf/transform.hpp b/cpp/include/cudf/transform.hpp
index e3830dd0a6cd..59a7becd5bb1 100644
--- a/cpp/include/cudf/transform.hpp
+++ b/cpp/include/cudf/transform.hpp
@@ -82,8 +82,8 @@ using transform_input = std::variant<column_view, scalar_column_view>;
  *
  *
  * @throws std::invalid_argument if any of the input columns have different sizes (except scalars)
- * @throws std::invalid_argument if `output_type` or any of the inputs are not fixed-width or string
- * types
+ * @throws std::invalid_argument if `output_type` or any of the inputs are not fixed-width, string,
+ * or dictionary types
  * @throws std::invalid_argument if any of the input columns have nulls
  * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and
  * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in
diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp
index 431ffa3340e1..cf67a5abddd0 100644
--- a/cpp/tests/transform/integration/unary_transform_test.cpp
+++ b/cpp/tests/transform/integration/unary_transform_test.cpp
@@ -333,8 +333,8 @@ __device__ inline void decode(cudf::string_view * output, cudf::string_view inpu
 
     cudf::transform_input inputs[] = {*a_encoded};
 
-    auto out =
-      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::STRING}, false);
+    auto out = cudf::transform_extended(
+      inputs, cuda, cudf::data_type{cudf::type_id::STRING}, cudf::udf_source_type::CUDA);
 
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
   }

From c654a31ea8b587050f2cb25d34f1f2fd2d82ce2f Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Fri, 13 Mar 2026 09:56:04 +0000
Subject: [PATCH 12/19] code review changes

---
 cpp/include/cudf/column/column_device_view_base.cuh    | 2 +-
 cpp/include/cudf/dictionary/dictionary_column_view.hpp | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index 0dea689de5ea..a462735afa3a 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -4,7 +4,7 @@
  */
 #pragma once
 
-#include <cudf/column/column_child_offsets.h>
+#include <cudf/column/column_child_offsets.hpp>
 #include <cudf/detail/offsets_iterator.cuh>
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/strings/string_view.cuh>
diff --git a/cpp/include/cudf/dictionary/dictionary_column_view.hpp b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
index 8666e0ec008d..dc82cfe0acf2 100644
--- a/cpp/include/cudf/dictionary/dictionary_column_view.hpp
+++ b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
@@ -4,7 +4,7 @@
  */
 #pragma once
 
-#include <cudf/column/column_child_offsets.h>
+#include <cudf/column/column_child_offsets.hpp>
 #include <cudf/column/column_view.hpp>
 
 /**

From 71c227c9ed122c876ec024c5f6f8658a1f77bd48 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Tue, 30 Jun 2026 08:44:41 +0000
Subject: [PATCH 13/19] rebase on latest cudf + code review changes

---
 .../cudf/column/column_child_offsets.hpp      |   7 +-
 .../cudf/column/column_device_view_base.cuh   |   6 +-
 .../dictionary/dictionary_column_view.hpp     |   2 +-
 cpp/include/cudf/transform.hpp                |  11 +-
 cpp/src/jit/column_accessor.cuh               |   7 +-
 cpp/src/jit/helpers.cpp                       |   2 +-
 cpp/src/transform/transform.cu                | 131 ++++++++++++++----
 .../integration/unary_transform_test.cpp      |  91 ++++++++++--
 8 files changed, 211 insertions(+), 46 deletions(-)

diff --git a/cpp/include/cudf/column/column_child_offsets.hpp b/cpp/include/cudf/column/column_child_offsets.hpp
index 0b9435f94338..594c00e5dd72 100644
--- a/cpp/include/cudf/column/column_child_offsets.hpp
+++ b/cpp/include/cudf/column/column_child_offsets.hpp
@@ -1,9 +1,14 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
 
+/**
+ * @file column_child_offsets.hpp
+ * @brief Constants for child column indices within compound column types
+ */
+
 #include <cudf/types.hpp>
 #include <cudf/utilities/export.hpp>
 
diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index e7376d853c4e..e48c91c7c05f 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -519,6 +519,7 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
     auto const& keys    = child(dictionary_keys_column_index);
     auto const index    = indices.template element<typename T::index_type>(element_index);
     return keys.template element<typename T::key_type>(index);
+  }
 
   /**
    * @brief Returns a nullable element at the specified index. If the element is null, returns
@@ -594,7 +595,8 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
 
 /**
  * @brief A mutable, non-owning view of device data as a column of elements
- * that is trivially copyable and usable in CUDA device code and offline-compiled code (i.e. NVRTC).
+ * that is trivially copyable and usable in CUDA device code and offline-compiled code (i.e.
+ * NVRTC).
  *
  * @ingroup column_classes
  */
diff --git a/cpp/include/cudf/dictionary/dictionary_column_view.hpp b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
index dc82cfe0acf2..df48faee484b 100644
--- a/cpp/include/cudf/dictionary/dictionary_column_view.hpp
+++ b/cpp/include/cudf/dictionary/dictionary_column_view.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
diff --git a/cpp/include/cudf/transform.hpp b/cpp/include/cudf/transform.hpp
index 5f1b2e6c1ba1..1d941223baa6 100644
--- a/cpp/include/cudf/transform.hpp
+++ b/cpp/include/cudf/transform.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -96,8 +96,9 @@ struct transform_output {
  *
  * @throws std::invalid_argument if any of the input columns have different sizes (except scalars)
  * @throws std::invalid_argument if `output_type` or any of the input types are not supported.
- * CUDA-supported types are fixed-width, string, and dictionary types, while PTX-supported types are integral and
- * floating-point types.
+ * CUDA-supported input types are fixed-width, string, and dictionary types. PTX-supported input
+ * types are fixed-width and dictionary types. CUDA-supported output types are fixed-width and
+ * string types. PTX-supported output types are fixed-width types.
  * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and
  * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in
  * this case.
@@ -145,6 +146,10 @@ std::unique_ptr<column> transform_extended(
  * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and
  * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in
  * this case.
+ * @throws std::invalid_argument if `output_type` or any of the input types are not supported.
+ * CUDA-supported input types are fixed-width, string, and dictionary types. PTX-supported input
+ * types are fixed-width and dictionary types. CUDA-supported output types are fixed-width and
+ * string types. PTX-supported output types are fixed-width types.
  * @throws std::invalid_argument if string offsets are provided for non-string output columns, or
  * if the number of string offsets does not match the number of output columns.
  *
diff --git a/cpp/src/jit/column_accessor.cuh b/cpp/src/jit/column_accessor.cuh
index ef1536cb409d..9851bbe4b868 100644
--- a/cpp/src/jit/column_accessor.cuh
+++ b/cpp/src/jit/column_accessor.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -45,7 +45,7 @@ struct column_accessor {
     return reinterpret_cast<column_type const&>(cols[index]);
   }
 
-  static __device__ element_type element(auto const* __restrict__ cols, size_type row)
+  static __device__ auto element(auto const* __restrict__ cols, size_type row)
   {
     return column(cols).template element<element_type>(map_index(row));
   }
@@ -60,8 +60,7 @@ struct column_accessor {
     return column(cols).is_valid(map_index(row));
   }
 
-  static __device__ optional_element_type nullable_element(auto const* __restrict__ cols,
-                                                           size_type row)
+  static __device__ auto nullable_element(auto const* __restrict__ cols, size_type row)
   {
     return column(cols).template nullable_element<element_type>(map_index(row));
   }
diff --git a/cpp/src/jit/helpers.cpp b/cpp/src/jit/helpers.cpp
index ed518a9f54d6..8056cb4d0532 100644
--- a/cpp/src/jit/helpers.cpp
+++ b/cpp/src/jit/helpers.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
diff --git a/cpp/src/transform/transform.cu b/cpp/src/transform/transform.cu
index 80b0509cfddf..117dc62922ab 100644
--- a/cpp/src/transform/transform.cu
+++ b/cpp/src/transform/transform.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -36,6 +36,10 @@
 namespace cudf {
 namespace {
 
+column_view as_column_view(scalar_column_view const& scalar) { return scalar.as_column_view(); }
+
+column_view as_column_view(column_view const& column) { return column; }
+
 struct mutable_fixed_width_column_view {
   mutable_column_view _view;
 
@@ -200,13 +204,53 @@ void launch(cudf::kernel const& kernel,
   kernel.launch({cfg.min_grid_size}, {cfg.block_size}, 0, stream, args);
 }
 
-std::string reflect_input_element(column_view const& c) { return type_to_name(c.type()); }
+namespace {
+
+std::string get_element_type_name(column_view const& view);
+
+struct element_type_name_fn {
+  template <typename T>
+  std::string operator()(column_view const& view) const
+    requires(is_fixed_width<T>() || std::same_as<T, cudf::string_view>)
+  {
+    return type_to_name(view.type());
+  }
 
-std::string reflect_input_element(scalar_column_view const& c) { return type_to_name(c.type()); }
+  template <typename T>
+  std::string operator()(column_view const& view) const
+    requires(std::same_as<T, cudf::dictionary32>)
+  {
+    return std::format("cudf::dictionary_element<{}, {}>",
+                       get_element_type_name(view.child(cudf::dictionary_indices_column_index)),
+                       get_element_type_name(view.child(cudf::dictionary_keys_column_index)));
+  }
+
+  template <typename T>
+  std::string operator()(column_view const& view) const
+    requires(!is_fixed_width<T>() && !std::same_as<T, cudf::string_view> &&
+             !std::same_as<T, cudf::dictionary32>)
+  {
+    CUDF_FAIL("Unsupported type for JIT compilation: " + type_to_name(view.type()));
+  }
+};
+
+std::string get_element_type_name(column_view const& view)
+{
+  return cudf::type_dispatcher(view.type(), element_type_name_fn{}, view);
+}
+
+}  // namespace
+
+std::string reflect_input_element(column_view const& c) { return get_element_type_name(c); }
+
+std::string reflect_input_element(scalar_column_view const& c)
+{
+  return get_element_type_name(c.as_column_view());
+}
 
 std::string reflect_output_element(fixed_width_column const& c)
 {
-  return type_to_name(c._col->type());
+  return get_element_type_name(c._col->view());
 }
 
 std::string reflect_output_element(string_views_column const&) { return "cudf::string_view"; }
@@ -216,6 +260,33 @@ std::string reflect_output_element(mutable_strings_column const&)
   return "cuda::std::span<char>";
 }
 
+std::string reflect_input_value_type(column_view const& c)
+{
+  return is_dictionary(c.type())
+           ? reflect_input_value_type(c.child(cudf::dictionary_keys_column_index))
+           : reflect_input_element(c);
+}
+
+std::string reflect_input_value_type(scalar_column_view const& c)
+{
+  return reflect_input_value_type(c.as_column_view());
+}
+
+std::string reflect_output_value_type(fixed_width_column const& c)
+{
+  return reflect_output_element(c);
+}
+
+std::string reflect_output_value_type(string_views_column const& c)
+{
+  return reflect_output_element(c);
+}
+
+std::string reflect_output_value_type(mutable_strings_column const& c)
+{
+  return reflect_output_element(c);
+}
+
 std::string reflect_input_column(column_view const&) { return "cudf::column_device_view_core"; }
 
 std::string reflect_input_column(scalar_column_view const&)
@@ -283,11 +354,12 @@ auto reflect(udf_source_type source_type,
 
   if (source_type == udf_source_type::PTX) {
     for (auto& in : inputs) {
-      ptx_in_types.push_back(std::visit([](auto& c) { return reflect_input_element(c); }, in));
+      ptx_in_types.push_back(std::visit([](auto& c) { return reflect_input_value_type(c); }, in));
     }
 
     for (auto& out : outputs) {
-      ptx_out_types.push_back(std::visit([](auto& c) { return reflect_output_element(c); }, out));
+      ptx_out_types.push_back(
+        std::visit([](auto& c) { return reflect_output_value_type(c); }, out));
     }
   }
 
@@ -511,17 +583,20 @@ void perform_checks(udf_source_type source_type,
                     std::span<std::unique_ptr<column> const> string_offsets)
 {
   if (source_type == udf_source_type::PTX) {
-    CUDF_EXPECTS(std::none_of(inputs.begin(),
-                              inputs.end(),
-                              [](auto& in) {
-                                return std::visit(
-                                  [](auto& c) {
-                                    return !is_integral(c.type()) && !is_floating_point(c.type());
-                                  },
-                                  in);
-                              }),
-                 "Transforms with PTX UDFs only support integer, floating-point, and boolean",
-                 std::invalid_argument);
+    static constexpr auto is_input_value_supported = [](auto const& c) {
+      return is_integral(c.type()) || is_floating_point(c.type());
+    };
+    static constexpr auto is_supported_input_type = [](auto const& c) {
+      auto col = std::visit([](auto& c) { return as_column_view(c); }, c);
+      return is_input_value_supported(col) ||
+             (is_dictionary(col.type()) &&
+              is_input_value_supported(col.child(dictionary_keys_column_index)));
+    };
+    CUDF_EXPECTS(
+      std::none_of(
+        inputs.begin(), inputs.end(), [](auto const& in) { return !is_supported_input_type(in); }),
+      "Transforms with PTX UDFs only support integer, floating-point, and boolean",
+      std::invalid_argument);
     CUDF_EXPECTS(std::none_of(outputs.begin(),
                               outputs.end(),
                               [](auto& out) {
@@ -542,14 +617,20 @@ void perform_checks(udf_source_type source_type,
                "Transforms only support output of fixed-width or string types",
                std::invalid_argument);
 
-  CUDF_EXPECTS(std::none_of(inputs.begin(),
-                            inputs.end(),
-                            [&](auto& in) {
-                              auto type = std::visit([](auto& c) { return c.type(); }, in);
-                              return !is_fixed_width(type) && type.id() != type_id::STRING;
-                            }),
-               "Transforms only support input of fixed-width or string types",
-               std::invalid_argument);
+  static constexpr auto is_input_value_supported = [](auto const& c) {
+    return is_fixed_width(c.type()) || c.type().id() == type_id::STRING || is_dictionary(c.type());
+  };
+  static constexpr auto is_supported_input_type = [&](auto const& c) {
+    auto col = std::visit([](auto const& c) { return as_column_view(c); }, c);
+    return is_input_value_supported(col) ||
+           (is_dictionary(col.type()) &&
+            is_input_value_supported(col.child(dictionary_keys_column_index)));
+  };
+  CUDF_EXPECTS(
+    std::none_of(
+      inputs.begin(), inputs.end(), [&](auto const& in) { return !is_supported_input_type(in); }),
+    "Transforms only support input of fixed-width, string, or dictionary types",
+    std::invalid_argument);
 
   if (!in_row_size.has_value()) {
     CUDF_EXPECTS(
diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp
index f5c71b8a47c3..7270a3f92e63 100644
--- a/cpp/tests/transform/integration/unary_transform_test.cpp
+++ b/cpp/tests/transform/integration/unary_transform_test.cpp
@@ -1,7 +1,7 @@
 /*
  * SPDX-FileCopyrightText: Copyright 2018-2019 BlazingDB, Inc.
  * SPDX-FileCopyrightText: Copyright 2018 Christian Noboa Mardini <christian@blazingdb.com>
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /*
@@ -414,7 +414,7 @@ __device__ inline void f(cudf::timestamp_us* output, cudf::timestamp_us input)
   test_udf<dtype>(cuda, op, data_init, 500, cudf::udf_source_type::CUDA);
 }
 
-TEST_F(UnaryOperationIntegrationTest, Transform_DictionaryString)
+TEST_F(UnaryOperationIntegrationTest, TransformDictionaryString)
 {
   std::string const cuda =
     R"***(
@@ -451,14 +451,14 @@ __device__ inline void decode(cudf::string_view * output, cudf::string_view inpu
 
     cudf::transform_input inputs[] = {*a_encoded};
 
-    auto out =
-      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::STRING}, false);
+    auto out = cudf::transform_extended(
+      inputs, cuda, cudf::data_type{cudf::type_id::STRING}, cudf::udf_source_type::CUDA);
 
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
   }
 }
 
-TEST_F(UnaryOperationIntegrationTest, Transform_DictionaryFloat)
+TEST_F(UnaryOperationIntegrationTest, TransformDictionaryFloat)
 {
   std::string const cuda =
     R"***(
@@ -466,6 +466,69 @@ __device__ inline void decode(float * output, float input){
   *output = input;
 })***";
 
+  // Generated from NUMBA, using:
+  //
+  // ```py
+  //
+  // from numba import cuda, float32
+  // from numba.cuda import compile_ptx_for_current_device
+  //
+  // # Define a CUDA device function
+  //
+  // @cuda.jit(device=True)
+  // def op(a):
+  //         return a
+  //
+  // # Define argument types for the function
+  // arg_types = (float32, )
+  //
+  // # Compile the device function as relocatable
+  // ptx, _ = cuda.compile_ptx_for_current_device(op, arg_types, device=True)
+  //
+  //
+  // # Print the PTX code
+  // print("Relocatable PTX Code:")
+  // print(ptx)
+  //
+  //
+  // ```
+  //
+  std::string const ptx = R"***(
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: UNKNOWN
+// Cuda compilation tools, release 13.2, V13.2.51
+// Based on NVVM 7.0.1
+//
+
+.version 9.2
+.target sm_86
+.address_size 64
+
+	// .globl	_ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf
+.common .global .align 8 .u64 _ZN08NumbaEnv8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf;
+
+.visible .func  (.param .b32 func_retval0) _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf(
+	.param .b64 _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_0,
+	.param .b32 _ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_1
+)
+{
+	.reg .f32 	%f<2>;
+	.reg .b32 	%r<2>;
+	.reg .b64 	%rd<2>;
+
+
+	ld.param.u64 	%rd1, [_ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_0];
+	ld.param.f32 	%f1, [_ZN8__main__3opxB2v1B96cw51cXTLSUwv1sCUt9Ww0FEw09RRQPKiLTj0gIGIFp_2b2oLQFEYYkHSQB1OQAk0Bynm21OizQ1K0UoIGvDpQE8oxrNQE_3dEf_param_1];
+	st.f32 	[%rd1], %f1;
+	mov.u32 	%r1, 0;
+	st.param.b32 	[func_retval0+0], %r1;
+	ret;
+
+}
+)***";
+
   // non-nullable
   {
     auto a = cudf::test::fixed_width_column_wrapper<float>(
@@ -474,10 +537,15 @@ __device__ inline void decode(float * output, float input){
     auto a_encoded                 = cudf::dictionary::encode(a->view());
     cudf::transform_input inputs[] = {*a_encoded};
 
-    auto out =
-      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, false);
+    auto out = cudf::transform_extended(
+      inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA);
 
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
+
+    auto out_ptx = cudf::transform_extended(
+      inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view());
   }
 
   // nullable
@@ -490,10 +558,15 @@ __device__ inline void decode(float * output, float input){
     auto a_encoded                 = cudf::dictionary::encode(a->view());
     cudf::transform_input inputs[] = {*a_encoded};
 
-    auto out =
-      cudf::transform_extended(inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, false);
+    auto out = cudf::transform_extended(
+      inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA);
 
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), a->view());
+
+    auto out_ptx = cudf::transform_extended(
+      inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view());
   }
 }
 

From 6c10d2e2031ff1f2f3db34cd142928135bb2251b Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Tue, 30 Jun 2026 09:46:46 +0100
Subject: [PATCH 14/19] Update cpp/include/cudf/column/column_child_offsets.hpp

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 cpp/include/cudf/column/column_child_offsets.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/include/cudf/column/column_child_offsets.hpp b/cpp/include/cudf/column/column_child_offsets.hpp
index 594c00e5dd72..89baa563e17e 100644
--- a/cpp/include/cudf/column/column_child_offsets.hpp
+++ b/cpp/include/cudf/column/column_child_offsets.hpp
@@ -17,7 +17,7 @@ namespace CUDF_EXPORT cudf {
 static constexpr size_type offsets_column_index = 0;  ///< Child index of the offsets column
 
 static constexpr size_type dictionary_indices_column_index =
-  0;  ///< Child index of the dictionary offsets column
+  0;  ///< Child index of the dictionary indices column
 
 static constexpr size_type dictionary_keys_column_index =
   1;  ///< Child index of the dictionary key column

From 98e9e3e23b657e3b09a3724d06865cdacd6c3f9f Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Tue, 30 Jun 2026 08:51:08 +0000
Subject: [PATCH 15/19] removed dead code for helper function

---
 cpp/src/jit/helpers.cpp | 12 ------------
 cpp/src/jit/helpers.hpp |  5 +----
 2 files changed, 1 insertion(+), 16 deletions(-)

diff --git a/cpp/src/jit/helpers.cpp b/cpp/src/jit/helpers.cpp
index 8056cb4d0532..a7620284b05c 100644
--- a/cpp/src/jit/helpers.cpp
+++ b/cpp/src/jit/helpers.cpp
@@ -81,18 +81,6 @@ std::map<uint32_t, std::string> build_ptx_params(std::span<std::string const> ou
   return params;
 }
 
-std::vector<std::string> input_type_names(
-  std::span<std::variant<column_view, scalar_column_view> const> views)
-{
-  std::vector<std::string> names;
-
-  std::transform(views.begin(), views.end(), std::back_inserter(names), [&](auto const& view) {
-    return std::visit([](auto& a) { return type_to_name(a.type()); }, view);
-  });
-
-  return names;
-}
-
 kernel get_udf_kernel(std::string const& source_file,
                       std::string const& kernel_name,
                       std::string const& cuda_source)
diff --git a/cpp/src/jit/helpers.hpp b/cpp/src/jit/helpers.hpp
index 8a7685a63ebd..9d06d7883c33 100644
--- a/cpp/src/jit/helpers.hpp
+++ b/cpp/src/jit/helpers.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -72,9 +72,6 @@ column_views_to_device(std::span<ColumnView const> views,
   return std::make_tuple(std::move(handles), std::move(device_array));
 }
 
-std::vector<std::string> input_type_names(
-  std::span<std::variant<column_view, scalar_column_view> const> views);
-
 kernel get_udf_kernel(std::string const& source_file,
                       std::string const& kernel_name,
                       std::string const& cuda_source);

From 374efecd58c032d3e8762096ebbe77174ef719a9 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Tue, 30 Jun 2026 09:32:24 +0000
Subject: [PATCH 16/19] code review changes

---
 .../cudf/column/column_device_view_base.cuh   |  3 ++-
 cpp/include/cudf/transform.hpp                | 18 ++++++++-------
 .../integration/unary_transform_test.cpp      | 22 +++++++++++++++++++
 3 files changed, 34 insertions(+), 9 deletions(-)

diff --git a/cpp/include/cudf/column/column_device_view_base.cuh b/cpp/include/cudf/column/column_device_view_base.cuh
index e48c91c7c05f..161aa969f4e0 100644
--- a/cpp/include/cudf/column/column_device_view_base.cuh
+++ b/cpp/include/cudf/column/column_device_view_base.cuh
@@ -517,7 +517,8 @@ class alignas(16) column_device_view_core : public detail::column_device_view_ba
   {
     auto const& indices = child(dictionary_indices_column_index);
     auto const& keys    = child(dictionary_keys_column_index);
-    auto const index    = indices.template element<typename T::index_type>(element_index);
+    auto const index    = indices.template element<typename T::index_type>(
+      element_index + offset());  // account for this view's _offset
     return keys.template element<typename T::key_type>(index);
   }
 
diff --git a/cpp/include/cudf/transform.hpp b/cpp/include/cudf/transform.hpp
index 1d941223baa6..c85a8549c9d1 100644
--- a/cpp/include/cudf/transform.hpp
+++ b/cpp/include/cudf/transform.hpp
@@ -95,10 +95,11 @@ struct transform_output {
  *
  *
  * @throws std::invalid_argument if any of the input columns have different sizes (except scalars)
- * @throws std::invalid_argument if `output_type` or any of the input types are not supported.
- * CUDA-supported input types are fixed-width, string, and dictionary types. PTX-supported input
- * types are fixed-width and dictionary types. CUDA-supported output types are fixed-width and
- * string types. PTX-supported output types are fixed-width types.
+ * @throws std::invalid_argument if any of the output or input types are not supported.
+ * CUDA-supported input types are fixed-width, string, and their dictionary types. PTX-supported
+ * input types are integrals, floats, and their dictionary types. CUDA-supported output types are
+ * fixed-width, string, and their dictionary types. PTX-supported output types are integrals,
+ * floats, and their dictionary types.
  * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and
  * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in
  * this case.
@@ -146,10 +147,11 @@ std::unique_ptr<column> transform_extended(
  * @throws std::invalid_argument if the inputs only have a scalar with no column inputs and
  * `row_size` is not provided. This is because the row size cannot be inferred from the inputs in
  * this case.
- * @throws std::invalid_argument if `output_type` or any of the input types are not supported.
- * CUDA-supported input types are fixed-width, string, and dictionary types. PTX-supported input
- * types are fixed-width and dictionary types. CUDA-supported output types are fixed-width and
- * string types. PTX-supported output types are fixed-width types.
+ * @throws std::invalid_argument if any of the output or input types are not supported.
+ * CUDA-supported input types are fixed-width, string, and their dictionary types. PTX-supported
+ * input types are integrals, floats, and their dictionary types. CUDA-supported output types are
+ * fixed-width, string, and their dictionary types. PTX-supported output types are integrals,
+ * floats, and their dictionary types.
  * @throws std::invalid_argument if string offsets are provided for non-string output columns, or
  * if the number of string offsets does not match the number of output columns.
  *
diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp
index 7270a3f92e63..033c658d217a 100644
--- a/cpp/tests/transform/integration/unary_transform_test.cpp
+++ b/cpp/tests/transform/integration/unary_transform_test.cpp
@@ -568,6 +568,28 @@ __device__ inline void decode(float * output, float input){
 
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view());
   }
+
+  // sliced
+  {
+    auto a = cudf::test::fixed_width_column_wrapper<float>(
+               {1.0F, 2.0F, 3.0F, 4.0F, 5.0F, 5.0F, 5.0F, 1.0F, 2.0F})
+               .release();
+
+    auto a_encoded     = cudf::dictionary::encode(a->view());
+    auto sliced_input  = cudf::slice(a_encoded->view(), {2, 7}).front();
+    auto sliced_expect = cudf::slice(a->view(), {2, 7}).front();
+
+    cudf::transform_input inputs[] = {sliced_input};
+
+    auto out = cudf::transform_extended(
+      inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA);
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out->view(), sliced_expect);
+
+    auto out_ptx = cudf::transform_extended(
+      inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX);
+
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view());
+  }
 }
 
 struct TernaryOperationTest : public cudf::test::BaseFixture {};

From 2ad4e8cd9bdfa69de25f8531a09bff0ccf8079b1 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Thu, 2 Jul 2026 15:32:44 +0000
Subject: [PATCH 17/19] Add test for empty column handling in
 transform_extended function

---
 .../transform/integration/unary_transform_test.cpp   | 12 ++++++++++++
 1 file changed, 12 insertions(+)

diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp
index 033c658d217a..d047d95985a3 100644
--- a/cpp/tests/transform/integration/unary_transform_test.cpp
+++ b/cpp/tests/transform/integration/unary_transform_test.cpp
@@ -569,6 +569,18 @@ __device__ inline void decode(float * output, float input){
     CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view());
   }
 
+  // empty column
+  {
+    auto a_empty                   = cudf::test::fixed_width_column_wrapper<float>({}).release();
+    auto a_encoded                 = cudf::dictionary::encode(a_empty->view());
+    cudf::transform_input inputs[] = {*a_encoded};
+
+    auto out = cudf::transform_extended(
+      inputs, cuda, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::CUDA);
+
+    EXPECT_EQ(out->size(), 0);
+  }
+
   // sliced
   {
     auto a = cudf::test::fixed_width_column_wrapper<float>(

From 66f1002a3ad135fae4257863ab2bcfe3c754ad3a Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Thu, 2 Jul 2026 16:22:00 +0000
Subject: [PATCH 18/19] Fix expected output in transform_extended test case

---
 cpp/tests/transform/integration/unary_transform_test.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/tests/transform/integration/unary_transform_test.cpp b/cpp/tests/transform/integration/unary_transform_test.cpp
index d047d95985a3..64738f54eb85 100644
--- a/cpp/tests/transform/integration/unary_transform_test.cpp
+++ b/cpp/tests/transform/integration/unary_transform_test.cpp
@@ -600,7 +600,7 @@ __device__ inline void decode(float * output, float input){
     auto out_ptx = cudf::transform_extended(
       inputs, ptx, cudf::data_type{cudf::type_id::FLOAT32}, cudf::udf_source_type::PTX);
 
-    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), a->view());
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(out_ptx->view(), sliced_expect);
   }
 }
 

From 0d92dda5c7947f1268c2107c9ac4d98d0f449854 Mon Sep 17 00:00:00 2001
From: Basit Ayantunde <rlamarrr@gmail.com>
Date: Fri, 3 Jul 2026 04:49:22 +0000
Subject: [PATCH 19/19] Add documentation for dictionary indices and keys
 column indices

---
 docs/cudf/source/libcudf/api_docs/dictionary_classes.rst | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst b/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst
index 00dec78c5f54..586c476afcb4 100644
--- a/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst
+++ b/docs/cudf/source/libcudf/api_docs/dictionary_classes.rst
@@ -1,5 +1,9 @@
 Dictionary Classes
 ==================
 
+.. doxygenvariable:: cudf::dictionary_indices_column_index
+
+.. doxygenvariable:: cudf::dictionary_keys_column_index
+
 .. doxygengroup:: dictionary_classes
    :members: