Support non-unique keys for dictionary column types

cpp/include/cudf/dictionary/dictionary_column_view.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -21,7 +21,8 @@ namespace CUDF_EXPORT cudf {
  * @brief A wrapper class for operations on a dictionary column.
  *
  * A dictionary column contains a set of keys and a column of indices.
- * The keys are a sorted set of unique values for the column.
+ * The keys may or may not be unique.
+ *
  * The indices represent the corresponding positions of each element's
  * value in the keys.
  */

cpp/include/cudf/dictionary/dictionary_factories.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
@@ -22,10 +22,7 @@ namespace CUDF_EXPORT cudf {
  * @brief Construct a dictionary column by copying the provided `keys`
  * and `indices`.
  *
- * It is expected that `keys_column.has_nulls() == false`.
- * It is assumed the elements in `keys_column` are unique and
- * are in a strict, total order. Meaning, `keys_column[i]` is ordered before
- * `keys_column[i+1]` for all `i in [0,n-1)` where `n` is the number of keys.
+ * The keys_column column must contain no nulls.
  *
  * The indices values must be in the range [0,keys_column.size()).
  *
@@ -45,7 +42,7 @@ namespace CUDF_EXPORT cudf {
  * @throw std::invalid_argument if keys_column contains nulls
  * @throw std::invalid_argument if indices_column type is not a signed integer
  *
- * @param keys_column Column of unique, ordered values to use as the new dictionary column's keys.
+ * @param keys_column Column of non-null values to use as the new dictionary column's keys.
  * @param indices_column Indices to use for the new dictionary column.
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned column's device memory.
@@ -62,16 +59,12 @@ std::unique_ptr<column> make_dictionary_column(
  * and indices columns.
  *
  * The keys_column and indices columns must contain no nulls.
- * It is assumed the elements in `keys_column` are unique and
- * are in a strict, total order. Meaning, `keys_column[i]` is ordered before
- * `keys_column[i+1]` for all `i in [0,n-1)` where `n` is the number of keys.
- *
  * The indices values must be in the range [0,keys_column.size()).
  *
  * @throw std::invalid_argument if keys_column or indices_column contains nulls
  * @throw std::invalid_argument if indices_column type is not a signed integer type
  *
- * @param keys_column Column of unique, ordered values to use as the new dictionary column's keys.
+ * @param keys_column Column of non-null values to use as the new dictionary column's keys.
  * @param indices_column Indices to use for the new dictionary column.
  * @param null_mask Null mask for the output column.
  * @param null_count Number of nulls for the output column.
@@ -97,7 +90,7 @@ std::unique_ptr<column> make_dictionary_column(std::unique_ptr<column> keys_colu
  *
  * @throw cudf::logic_error if keys_column contains nulls
  *
- * @param keys_column Column of unique, ordered values to use as the new dictionary column's keys.
+ * @param keys_column Column of non-null values to use as the new dictionary column's keys.
  * @param indices_column Indices values and null-mask to use for the new dictionary column.
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned column's device memory.

cpp/include/cudf/dictionary/update_keys.hpp

@@ -117,6 +117,9 @@ std::unique_ptr<column> remove_unused_keys(
  * The indices are updated to reflect the position values of the new keys.
  * Any indices pointing to removed keys sets those rows to null.
  *
+ * Although duplicate keys are allowed, indices in the returned dictionary may
+ * only reference one of the duplicates.
+ *
  * @code{.pseudo}
  * d1 = {keys=["a", "b", "c"], indices=[2, 0, 1, 2, 1]}
  * d2 = set_keys(existing_dict, ["b","c","d"])

cpp/src/reductions/argmax.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -15,7 +15,7 @@ std::unique_ptr<scalar> argmax(column_view const& input,
                                rmm::device_async_resource_ref mr)
 {
   auto const dispatch_type =
-    is_dictionary(input.type()) ? dictionary_column_view(input).indices().type() : input.type();
+    is_dictionary(input.type()) ? dictionary_column_view(input).keys().type() : input.type();
   return type_dispatcher(
     dispatch_type, simple::detail::arg_minmax_dispatcher<aggregation::ARGMAX>{}, input, stream, mr);
 }

cpp/src/reductions/argmin.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -15,7 +15,7 @@ std::unique_ptr<scalar> argmin(column_view const& input,
                                rmm::device_async_resource_ref mr)
 {
   auto const dispatch_type =
-    is_dictionary(input.type()) ? dictionary_column_view(input).indices().type() : input.type();
+    is_dictionary(input.type()) ? dictionary_column_view(input).keys().type() : input.type();
   return type_dispatcher(
     dispatch_type, simple::detail::arg_minmax_dispatcher<aggregation::ARGMIN>{}, input, stream, mr);
 }

cpp/src/reductions/extrema_utils.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
  */
 
@@ -8,6 +8,7 @@
 #include "nested_types_extrema_utils.cuh"
 
 #include <cudf/aggregation.hpp>
+#include <cudf/dictionary/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/scalar/scalar_factories.hpp>
 #include <cudf/utilities/traits.hpp>
@@ -98,6 +99,9 @@ class arg_minmax_dispatcher {
   }
 
   // This function is used for types such as string, timestamp, fixed point, etc.
+  // `input` may be the dictionary column itself (rather than its keys or a decoded column):
+  // the lexicographic self_comparator compares `keys[indices[i]]` for dictionary columns
+  // directly, so no decoding is required and no assumption is made that the keys are sorted.
   template <typename ElementType>
   [[nodiscard]] size_type find_arg_minmax(column_view const& input,
                                           rmm::cuda_stream_view stream) const
@@ -121,10 +125,26 @@ class arg_minmax_dispatcher {
                                           rmm::cuda_stream_view stream) const
     requires(cudf::is_numeric<ElementType>())  // integer + floating point numbers
   {
+    using Op = std::conditional_t<K == aggregation::ARGMIN,
+                                  reduction::detail::op::min,
+                                  reduction::detail::op::max>;
+    // Dictionary keys are not guaranteed to be sorted, so the index of the min/max index does
+    // not identify the min/max key. Read `keys[indices[i]]` directly per row via a lazy
+    // iterator instead of decoding (and copying) the whole column.
+    if (is_dictionary(input.type())) {
+      auto const d_dict = column_device_view::create(input, stream);
+      if (input.has_nulls()) {
+        auto const transformer =
+          Op{}.template get_null_replacing_element_transformer<ElementType>();
+        auto const p =
+          cudf::dictionary::detail::make_dictionary_pair_iterator<ElementType>(*d_dict, true);
+        auto const it = thrust::make_transform_iterator(p, transformer);
+        return find_extremum_idx(it, input.size(), stream);
+      }
+      auto const it = cudf::dictionary::detail::make_dictionary_iterator<ElementType>(*d_dict);
+      return find_extremum_idx(it, input.size(), stream);
+    }
     if (input.has_nulls()) {
-      using Op               = std::conditional_t<K == aggregation::ARGMIN,
-                                                  reduction::detail::op::min,
-                                                  reduction::detail::op::max>;
       auto const d_input     = column_device_view::create(input, stream);
       auto const transformer = Op{}.template get_null_replacing_element_transformer<ElementType>();
       auto const it =
@@ -150,9 +170,7 @@ class arg_minmax_dispatcher {
                                                    rmm::device_async_resource_ref mr) const
     requires(is_supported<ElementType>())
   {
-    auto const& values =
-      is_dictionary(input.type()) ? dictionary_column_view(input).get_indices_annotated() : input;
-    auto const idx = find_arg_minmax<ElementType>(values, stream);
+    auto const idx = find_arg_minmax<ElementType>(input, stream);
     return make_fixed_width_scalar<size_type>(idx, stream, mr);
   }
 

cpp/src/reductions/minmax.cu

@@ -1,16 +1,16 @@
 /*
- * 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
  */
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_view.hpp>
-#include <cudf/detail/copy.hpp>
 #include <cudf/detail/device_scalar.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/device_operators.cuh>
+#include <cudf/dictionary/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/reduction.hpp>
 #include <cudf/reduction/detail/reduction_functions.hpp>
@@ -129,6 +129,94 @@ struct create_minmax_with_nulls {
   }
 };
 
+/**
+ * @brief Functor to copy a minmax_pair result to individual scalar instances.
+ *
+ * @tparam T type of the data
+ * @tparam ResultType result type to assign min, max to minmax_pair<T>
+ */
+template <typename T, typename ResultType = minmax_pair<T>>
+struct assign_min_max {
+  __device__ void operator()()
+  {
+    *min_data = result->min_val;
+    *max_data = result->max_val;
+  }
+
+  ResultType* result;
+  T* min_data;
+  T* max_data;
+};
+
+/**
+ * @brief Computes a minmax_pair<T> reduction directly over a dictionary column's decoded key
+ * values, i.e. `keys[indices[i]]` for each row `i`.
+ *
+ * @tparam T The dictionary's key type
+ */
+template <typename T>
+auto reduce_dictionary(column_view const& col, rmm::cuda_stream_view stream)
+{
+  auto d_dictionary = column_device_view::create(col, stream);
+  if (col.has_nulls()) {
+    auto pair_to_minmax = thrust::make_transform_iterator(
+      cudf::dictionary::detail::make_dictionary_pair_iterator<T>(*d_dictionary, true),
+      create_minmax_with_nulls<T>{});
+    return reduce_device(pair_to_minmax, col.size(), minmax_binary_op<T>{}, stream);
+  } else {
+    auto col_to_minmax = thrust::make_transform_iterator(
+      cudf::dictionary::detail::make_dictionary_iterator<T>(*d_dictionary), create_minmax<T>{});
+    return reduce_device(col_to_minmax, col.size(), minmax_binary_op<T>{}, stream);
+  }
+}
+
+/**
+ * @brief Dispatch functor for minmax operation on a dictionary column, dispatched on the
+ * dictionary's key type.
+ */
+struct minmax_dictionary_functor {
+  template <typename T>
+  static constexpr bool is_supported()
+  {
+    return !cudf::is_dictionary<T>() && !std::is_same_v<T, cudf::list_view> &&
+           !std::is_same_v<T, cudf::struct_view>;
+  }
+
+  template <typename T>
+  std::pair<std::unique_ptr<scalar>, std::unique_ptr<scalar>> operator()(
+    column_view const& col, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
+    requires(is_supported<T>() and !std::is_same_v<T, cudf::string_view>)
+  {
+    using storage_type = device_storage_type_t<T>;
+    auto dev_result    = reduce_dictionary<storage_type>(col, stream);
+    using ScalarType   = cudf::scalar_type_t<T>;
+    auto minimum       = std::make_unique<ScalarType>(T{}, true, stream, mr);
+    auto maximum       = std::make_unique<ScalarType>(T{}, true, stream, mr);
+    cudf::detail::device_single_thread(
+      assign_min_max<storage_type>{dev_result.data(), minimum->data(), maximum->data()}, stream);
+    return {std::move(minimum), std::move(maximum)};
+  }
+
+  template <typename T>
+  std::pair<std::unique_ptr<scalar>, std::unique_ptr<scalar>> operator()(
+    column_view const& col, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
+    requires(std::is_same_v<T, cudf::string_view>)
+  {
+    auto dev_result        = reduce_dictionary<cudf::string_view>(col, stream);
+    auto const host_result = dev_result.value(stream);
+    return {std::make_unique<string_scalar>(host_result.min_val, true, stream, mr),
+            std::make_unique<string_scalar>(host_result.max_val, true, stream, mr)};
+  }
+
+  template <typename T>
+  std::pair<std::unique_ptr<scalar>, std::unique_ptr<scalar>> operator()(
+    column_view const&, rmm::cuda_stream_view, rmm::device_async_resource_ref)
+    requires(!is_supported<T>())
+  {
+    CUDF_FAIL("dictionary key type not supported for minmax() operation");
+  }
+};
+
 /**
  * @brief Dispatch functor for minmax operation.
  *
@@ -160,25 +248,6 @@ struct minmax_functor {
     }
   }
 
-  /**
-   * @brief Functor to copy a minmax_pair result to individual scalar instances.
-   *
-   * @tparam T type of the data
-   * @tparam ResultType result type to assign min, max to minmax_pair<T>
-   */
-  template <typename T, typename ResultType = minmax_pair<T>>
-  struct assign_min_max {
-    __device__ void operator()()
-    {
-      *min_data = result->min_val;
-      *max_data = result->max_val;
-    }
-
-    ResultType* result;
-    T* min_data;
-    T* max_data;
-  };
-
   template <typename T>
   std::pair<std::unique_ptr<scalar>, std::unique_ptr<scalar>> operator()(
     cudf::column_view const& col, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
@@ -223,15 +292,8 @@ struct minmax_functor {
     cudf::column_view const& col, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
     requires(cudf::is_dictionary<T>())
   {
-    // computes minimum and maximum on the dictionary indices as dictionary32 values
-    auto d_indices     = reduce<T>(col, stream);
-    auto const indices = d_indices.value(stream);
-    // use these values to slice the keys column (add 1 for complete inclusion)
-    auto keys = cudf::detail::slice(dictionary_column_view(col).keys(),
-                                    {indices.min_val.value(), indices.max_val.value() + 1},
-                                    stream)
-                  .front();
-    return type_dispatcher(keys.type(), minmax_functor{}, keys, stream, mr);
+    auto const keys_type = dictionary_column_view(col).keys().type();
+    return type_dispatcher(keys_type, minmax_dictionary_functor{}, col, stream, mr);
   }
 
   template <typename T>

cpp/tests/dictionary/add_keys_test.cpp

@@ -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
  */
 
@@ -58,6 +58,29 @@ TEST_F(DictionaryAddKeysTest, WithNull)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(decoded->view(), input);  // new keys should not change anything
 }
 
+TEST_F(DictionaryAddKeysTest, DuplicateKeys)
+{
+  auto input = cudf::test::dictionary_column_wrapper<std::string>(
+    {"eee", "aaa", "ddd", "bbb", "ccc", "ccc", "ccc", "eee", "aaa"});
+
+  auto const dup_keys  = cudf::test::strings_column_wrapper{"aaa", "ccc", "eee", "ccc"};
+  auto const with_dups = cudf::dictionary::set_keys(input, dup_keys);  // force duplicate keys
+
+  // add_keys with a key already present
+  auto const added_keys  = cudf::test::strings_column_wrapper{"ccc", "fff"};
+  auto const result      = cudf::dictionary::add_keys(with_dups->view(), added_keys);
+  auto const result_view = cudf::dictionary_column_view(result->view());
+
+  // old 4 keys + "fff" only — "ccc" must not have been added a third time
+  EXPECT_EQ(result_view.keys_size(), 5);
+
+  // indices are unchanged so decode should match what set_keys produced
+  auto expected_decoded = cudf::test::strings_column_wrapper(
+    {"eee", "aaa", "", "", "ccc", "ccc", "ccc", "eee", "aaa"}, {1, 1, 0, 0, 1, 1, 1, 1, 1});
+  auto const decoded = cudf::dictionary::decode(result_view);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(decoded->view(), expected_decoded);
+}
+
 TEST_F(DictionaryAddKeysTest, Errors)
 {
   cudf::test::fixed_width_column_wrapper<int64_t> input{1, 2, 3};