diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 12b08b136c6..b3ad7cda1b7 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -795,8 +795,6 @@ add_library(
   src/structs/structs_column_factories.cu
   src/structs/structs_column_view.cpp
   src/structs/utilities.cu
-  src/table/primitive_row_operators.cu
-  src/table/row_operators.cu
   src/table/table.cpp
   src/table/table_device_view.cu
   src/table/table_view.cpp
diff --git a/cpp/include/cudf/detail/row_operator/row_operators.cuh b/cpp/include/cudf/detail/row_operator/row_operators.cuh
index 9033bed5bec..1209f1bdccb 100644
--- a/cpp/include/cudf/detail/row_operator/row_operators.cuh
+++ b/cpp/include/cudf/detail/row_operator/row_operators.cuh
@@ -28,7 +28,6 @@
 #include <cudf/lists/lists_column_device_view.cuh>
 #include <cudf/sorting.hpp>
 #include <cudf/structs/structs_column_device_view.cuh>
-#include <cudf/table/row_operators.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/traits.hpp>
diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
deleted file mode 100644
index e2e08c59f55..00000000000
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ /dev/null
@@ -1,2111 +0,0 @@
-/*
- * Copyright (c) 2022-2025, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-/**
- * @file
- * @deprecated This header is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-
-#include <cudf/column/column_device_view.cuh>
-#include <cudf/detail/iterator.cuh>
-#include <cudf/detail/utilities/algorithm.cuh>
-#include <cudf/detail/utilities/assert.cuh>
-#include <cudf/hashing/detail/default_hash.cuh>
-#include <cudf/hashing/detail/hashing.hpp>
-#include <cudf/lists/detail/dremel.hpp>
-#include <cudf/lists/list_device_view.cuh>
-#include <cudf/lists/lists_column_device_view.cuh>
-#include <cudf/sorting.hpp>
-#include <cudf/structs/structs_column_device_view.cuh>
-#include <cudf/table/row_operators.cuh>
-#include <cudf/table/table_device_view.cuh>
-#include <cudf/utilities/default_stream.hpp>
-#include <cudf/utilities/traits.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
-
-#include <cuda/std/functional>
-#include <cuda/std/limits>
-#include <cuda/std/optional>
-#include <cuda/std/tuple>
-#include <cuda/std/utility>
-#include <thrust/detail/use_default.h>
-#include <thrust/equal.h>
-#include <thrust/execution_policy.h>
-#include <thrust/functional.h>
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/iterator_adaptor.h>
-#include <thrust/iterator/iterator_categories.h>
-#include <thrust/iterator/iterator_facade.h>
-#include <thrust/iterator/transform_iterator.h>
-#include <thrust/logical.h>
-#include <thrust/swap.h>
-#include <thrust/transform_reduce.h>
-
-#include <memory>
-#include <type_traits>
-
-namespace CUDF_EXPORT cudf {
-
-namespace row::primitive {
-class row_equality_comparator;  // Forward declaration
-
-template <template <typename> class Hash>
-class row_hasher;  // Forward declaration
-}  // namespace row::primitive
-
-namespace experimental {
-
-/**
- * @brief A map from cudf::type_id to cudf type that excludes LIST and STRUCT types.
- *
- * To be used with type_dispatcher in place of the default map, when it is required that STRUCT and
- * LIST map to void. This is useful when we want to avoid recursion in a functor. For example, in
- * element_comparator, we have a specialization for STRUCT but the type_dispatcher in it is only
- * used to dispatch to the same functor for non-nested types. Even when we're guaranteed to not have
- * non-nested types at that point, the compiler doesn't know this and would try to create recursive
- * code which is very slow.
- *
- * Usage:
- * @code
- * type_dispatcher<dispatch_nested_to_void>(data_type(), functor{});
- * @endcode
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <cudf::type_id t>
-struct dispatch_void_if_nested {
-  /// The type to dispatch to if the type is nested
-  using type = std::conditional_t<t == type_id::STRUCT or t == type_id::LIST, void, id_to_type<t>>;
-};
-
-namespace row {
-
-/// Strongly typed index for left-hand side table rows
-enum class lhs_index_type : size_type {};
-/// Strongly typed index for right-hand side table rows
-enum class rhs_index_type : size_type {};
-
-/**
- * @brief A counting iterator that uses strongly typed indices bound to tables.
- *
- * Performing lexicographic or equality comparisons between values in two
- * tables requires the use of strongly typed indices. The strong index types
- * `lhs_index_type` and `rhs_index_type` ensure that index values are bound to
- * the correct table, regardless of the order in which these indices are
- * provided to the call operator. This struct and its type aliases
- * `lhs_iterator` and `rhs_iterator` provide an interface similar to a counting
- * iterator, with strongly typed values to represent the table indices.
- *
- * @tparam Index The strong index type
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <typename Index, typename Underlying = std::underlying_type_t<Index>>
-struct strong_index_iterator : public thrust::iterator_facade<strong_index_iterator<Index>,
-                                                              Index,
-                                                              thrust::use_default,
-                                                              thrust::random_access_traversal_tag,
-                                                              Index,
-                                                              Underlying> {
-  using super_t =
-    thrust::iterator_adaptor<strong_index_iterator<Index>, Index>;  ///< The base class
-
-  /**
-   * @brief Constructs a strong index iterator
-   *
-   * @param n The beginning index
-   */
-  explicit constexpr strong_index_iterator(Underlying n) : begin{n} {}
-
-  friend class thrust::iterator_core_access;  ///< Allow access to the base class
-
- private:
-  __device__ constexpr void increment() { ++begin; }
-  __device__ constexpr void decrement() { --begin; }
-
-  __device__ constexpr void advance(Underlying n) { begin += n; }
-
-  __device__ constexpr bool equal(strong_index_iterator<Index> const& other) const noexcept
-  {
-    return begin == other.begin;
-  }
-
-  __device__ constexpr Index dereference() const noexcept { return static_cast<Index>(begin); }
-
-  __device__ constexpr Underlying distance_to(
-    strong_index_iterator<Index> const& other) const noexcept
-  {
-    return other.begin - begin;
-  }
-
-  Underlying begin{};
-};
-
-/**
- * @brief Iterator representing indices into a left-side table.
- */
-using lhs_iterator = strong_index_iterator<lhs_index_type>;
-
-/**
- * @brief Iterator representing indices into a right-side table.
- */
-using rhs_iterator = strong_index_iterator<rhs_index_type>;
-
-namespace lexicographic {
-
-/**
- * @brief Computes a weak ordering of two values with special sorting behavior.
- *
- * This relational comparator functor compares physical values rather than logical
- * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or
- * greater than other values and is IEEE-754 compliant.
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-struct physical_element_comparator {
-  /**
-   * @brief Operator for relational comparisons.
-   *
-   * @param lhs First element
-   * @param rhs Second element
-   * @return Relation between elements
-   */
-  template <typename Element>
-  __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept
-  {
-    return detail::compare_elements(lhs, rhs);
-  }
-};
-
-/**
- * @brief Relational comparator functor that compares physical values rather than logical
- * elements like lists, strings, or structs. It evaluates `NaN` as equivalent to other `NaN`s and
- * greater than all other values.
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-struct sorting_physical_element_comparator {
-  /**
-   * @brief Operator for relational comparison of non-floating point values.
-   *
-   * @param lhs First element
-   * @param rhs Second element
-   * @return Relation between elements
-   */
-  template <typename Element, CUDF_ENABLE_IF(not std::is_floating_point_v<Element>)>
-  __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept
-  {
-    return detail::compare_elements(lhs, rhs);
-  }
-
-  /**
-   * @brief Operator for relational comparison of floating point values.
-   *
-   * @param lhs First element
-   * @param rhs Second element
-   * @return Relation between elements
-   */
-  template <typename Element, CUDF_ENABLE_IF(std::is_floating_point_v<Element>)>
-  __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept
-  {
-    if (isnan(lhs)) {
-      return isnan(rhs) ? weak_ordering::EQUIVALENT : weak_ordering::GREATER;
-    } else if (isnan(rhs)) {
-      return weak_ordering::LESS;
-    }
-
-    return detail::compare_elements(lhs, rhs);
-  }
-};
-
-/// Optional dremel device view for handling nested column structures
-using optional_dremel_view = cuda::std::optional<detail::dremel_device_view const>;
-
-// The has_nested_columns template parameter of the device_row_comparator is
-// necessary to help the compiler optimize our code. Without it, the list and
-// struct view specializations are present in the code paths used for primitive
-// types, and the compiler fails to inline this nearly as well resulting in a
-// significant performance drop.  As a result, there is some minor tension in
-// the current design between the presence of this parameter and the way that
-// the Dremel data is passed around, first as a
-// std::optional<device_span<dremel_device_view>> in the
-// preprocessed_table/device_row_comparator (which is always valid when
-// has_nested_columns and is otherwise invalid) that is then unpacked to a
-// cuda::std::optional<dremel_device_view> at the element_comparator level (which
-// is always valid for a list column and otherwise invalid).  We cannot use an
-// additional template parameter for the element_comparator on a per-column
-// basis because we cannot conditionally define dremel_device_view member
-// variables without jumping through extra hoops with inheritance, so the
-// cuda::std::optional<dremel_device_view> member must be an optional rather than
-// a raw dremel_device_view.
-/**
- * @brief Computes the lexicographic comparison between 2 rows.
- *
- * Lexicographic ordering is determined by:
- * - Two rows are compared element by element.
- * - The first mismatching element defines which row is lexicographically less
- * or greater than the other.
- * - If the rows are compared without mismatched elements, the rows are equivalent
- *
- *
- * Lexicographic ordering is exactly equivalent to doing an alphabetical sort of
- * two words, for example, `aac` would be *less* than (or precede) `abb`. The
- * second letter in both words is the first non-equal letter, and `a < b`, thus
- * `aac < abb`.
- *
- * @note The operator overloads in sub-class `element_comparator` are templated via the
- *        `type_dispatcher` to help select an overload instance for each column in a table.
- *        So, `cudf::is_nested<Element>` will return `true` if the table has nested-type columns,
- *        but it will be a runtime error if template parameter `has_nested_columns != true`.
- *
- * @tparam has_nested_columns compile-time optimization for primitive types.
- *         This template parameter is to be used by the developer by querying
- *         `cudf::has_nested_columns(input)`. `true` compiles operator
- *         overloads for nested types, while `false` only compiles operator
- *         overloads for primitive types.
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- * @tparam PhysicalElementComparator A relational comparator functor that compares individual values
- * rather than logical elements, defaults to `NaN` aware relational comparator that evaluates `NaN`
- * as greater than all other values.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <bool has_nested_columns,
-          typename Nullate,
-          typename PhysicalElementComparator = sorting_physical_element_comparator>
-class device_row_comparator {
- public:
-  friend class self_comparator;       ///< Allow self_comparator to access private members
-  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
-
-  /**
-   * @brief Construct a function object for performing a lexicographic
-   * comparison between the rows of two tables.
-   *
-   * @param check_nulls Indicates if any input column contains nulls.
-   * @param lhs The first table
-   * @param rhs The second table (may be the same table as `lhs`)
-   * @param l_dremel_device_views lhs table dremel device view for list type
-   * @param r_dremel_device_views rhs table dremel device view for list type
-   * @param depth Optional, device array the same length as a row that contains starting depths of
-   * columns if they're nested, and 0 otherwise.
-   * @param column_order Optional, device array the same length as a row that indicates the desired
-   * ascending/descending order of each column in a row. If `nullopt`, it is assumed all columns are
-   * sorted in ascending order.
-   * @param null_precedence Optional, device array the same length as a row and indicates how null
-   * values compare to all other for every column. If `nullopt`, then null precedence would be
-   * `null_order::BEFORE` for all columns.
-   * @param comparator Physical element relational comparison functor.
-   */
-  device_row_comparator(
-    Nullate check_nulls,
-    table_device_view lhs,
-    table_device_view rhs,
-    device_span<detail::dremel_device_view const> l_dremel_device_views,
-    device_span<detail::dremel_device_view const> r_dremel_device_views,
-    cuda::std::optional<device_span<int const>> depth                  = cuda::std::nullopt,
-    cuda::std::optional<device_span<order const>> column_order         = cuda::std::nullopt,
-    cuda::std::optional<device_span<null_order const>> null_precedence = cuda::std::nullopt,
-    PhysicalElementComparator comparator                               = {}) noexcept
-    : _lhs{lhs},
-      _rhs{rhs},
-      _l_dremel(l_dremel_device_views),
-      _r_dremel(r_dremel_device_views),
-      _check_nulls{check_nulls},
-      _depth{depth},
-      _column_order{column_order},
-      _null_precedence{null_precedence},
-      _comparator{comparator}
-  {
-  }
-
-  /**
-   * @brief Construct a function object for performing a lexicographic
-   * comparison between the rows of two tables.
-   * This is a special overload to allow device-side construction of the
-   * comparator for cases where no preprocessing is needed, i.e. tables with
-   * non-nested type columns.
-   *
-   * @param check_nulls Indicates if any input column contains nulls.
-   * @param lhs The first table
-   * @param rhs The second table (may be the same table as `lhs`)
-   * @param column_order Optional, device array the same length as a row that indicates the desired
-   * ascending/descending order of each column in a row. If `nullopt`, it is assumed all columns are
-   * sorted in ascending order.
-   * @param null_precedence Optional, device array the same length as a row and indicates how null
-   * values compare to all other for every column. If `nullopt`, then null precedence would be
-   * `null_order::BEFORE` for all columns.
-   * @param comparator Physical element relational comparison functor.
-   */
-  template <bool nested_disable = not has_nested_columns, CUDF_ENABLE_IF(nested_disable)>
-  __device__ device_row_comparator(
-    Nullate check_nulls,
-    table_device_view lhs,
-    table_device_view rhs,
-    cuda::std::optional<device_span<order const>> column_order         = cuda::std::nullopt,
-    cuda::std::optional<device_span<null_order const>> null_precedence = cuda::std::nullopt,
-    PhysicalElementComparator comparator                               = {}) noexcept
-    : _lhs{lhs},
-      _rhs{rhs},
-      _l_dremel{},
-      _r_dremel{},
-      _check_nulls{check_nulls},
-      _depth{},
-      _column_order{column_order},
-      _null_precedence{null_precedence},
-      _comparator{comparator}
-  {
-  }
-
-  /**
-   * @brief Performs a relational comparison between two elements in two columns.
-   */
-  class element_comparator {
-   public:
-    /**
-     * @brief Construct type-dispatched function object for performing a
-     * relational comparison between two elements.
-     *
-     * @note `lhs` and `rhs` may be the same.
-     *
-     * @param check_nulls Indicates if either input column contains nulls.
-     * @param lhs The column containing the first element
-     * @param rhs The column containing the second element (may be the same as lhs)
-     * @param null_precedence Indicates how null values are ordered with other values
-     * @param depth The depth of the column if part of a nested column @see
-     * preprocessed_table::depths
-     * @param comparator Physical element relational comparison functor.
-     * @param l_dremel_device_view <>
-     * @param r_dremel_device_view <>
-     */
-    __device__ element_comparator(Nullate check_nulls,
-                                  column_device_view lhs,
-                                  column_device_view rhs,
-                                  null_order null_precedence                = null_order::BEFORE,
-                                  int depth                                 = 0,
-                                  PhysicalElementComparator comparator      = {},
-                                  optional_dremel_view l_dremel_device_view = {},
-                                  optional_dremel_view r_dremel_device_view = {})
-      : _lhs{lhs},
-        _rhs{rhs},
-        _check_nulls{check_nulls},
-        _null_precedence{null_precedence},
-        _depth{depth},
-        _l_dremel_device_view{l_dremel_device_view},
-        _r_dremel_device_view{r_dremel_device_view},
-        _comparator{comparator}
-    {
-    }
-
-    /**
-     * @brief Performs a relational comparison between the specified elements
-     *
-     * @param lhs_element_index The index of the first element
-     * @param rhs_element_index The index of the second element
-     * @return Indicates the relationship between the elements in the `lhs` and `rhs` columns, along
-     * with the depth at which a null value was encountered.
-     */
-    template <typename Element,
-              CUDF_ENABLE_IF(cudf::is_relationally_comparable<Element, Element>())>
-    __device__ cuda::std::pair<weak_ordering, int> operator()(
-      size_type const lhs_element_index, size_type const rhs_element_index) const noexcept
-    {
-      if (_check_nulls) {
-        bool const lhs_is_null{_lhs.is_null(lhs_element_index)};
-        bool const rhs_is_null{_rhs.is_null(rhs_element_index)};
-
-        if (lhs_is_null or rhs_is_null) {  // at least one is null
-          return cuda::std::pair(null_compare(lhs_is_null, rhs_is_null, _null_precedence), _depth);
-        }
-      }
-
-      return cuda::std::pair(_comparator(_lhs.element<Element>(lhs_element_index),
-                                         _rhs.element<Element>(rhs_element_index)),
-                             cuda::std::numeric_limits<int>::max());
-    }
-
-    /**
-     * @brief Throws run-time error when columns types cannot be compared
-     *        or if this class is instantiated with `has_nested_columns = false` but
-     *        passed tables with nested columns
-     *
-     * @return Ordering
-     */
-    template <typename Element,
-              CUDF_ENABLE_IF(not cudf::is_relationally_comparable<Element, Element>() and
-                             (not has_nested_columns or not cudf::is_nested<Element>()))>
-    __device__ cuda::std::pair<weak_ordering, int> operator()(size_type const,
-                                                              size_type const) const noexcept
-    {
-      CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
-    }
-
-    /**
-     * @brief Compares two struct-type columns
-     *
-     * @param lhs_element_index The index of the first element
-     * @param rhs_element_index The index of the second element
-     * @return Indicates the relationship between the elements in the `lhs` and `rhs` columns, along
-     * with the depth at which a null value was encountered.
-     */
-    template <typename Element,
-              CUDF_ENABLE_IF(has_nested_columns and std::is_same_v<Element, cudf::struct_view>)>
-    __device__ cuda::std::pair<weak_ordering, int> operator()(
-      size_type const lhs_element_index, size_type const rhs_element_index) const noexcept
-    {
-      column_device_view lcol = _lhs;
-      column_device_view rcol = _rhs;
-      int depth               = _depth;
-      while (lcol.type().id() == type_id::STRUCT) {
-        bool const lhs_is_null{lcol.is_null(lhs_element_index)};
-        bool const rhs_is_null{rcol.is_null(rhs_element_index)};
-
-        if (lhs_is_null or rhs_is_null) {  // at least one is null
-          weak_ordering state = null_compare(lhs_is_null, rhs_is_null, _null_precedence);
-          return cuda::std::pair(state, depth);
-        }
-
-        if (lcol.num_child_columns() == 0) {
-          return cuda::std::pair(weak_ordering::EQUIVALENT, cuda::std::numeric_limits<int>::max());
-        }
-
-        // Non-empty structs have been modified to only have 1 child when using this.
-        lcol = detail::structs_column_device_view(lcol).get_sliced_child(0);
-        rcol = detail::structs_column_device_view(rcol).get_sliced_child(0);
-        ++depth;
-      }
-
-      return cudf::type_dispatcher<dispatch_void_if_nested>(
-        lcol.type(),
-        element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth, _comparator},
-        lhs_element_index,
-        rhs_element_index);
-    }
-
-    /**
-     * @brief Compares two list-type columns
-     *
-     * @param lhs_element_index The index of the first element
-     * @param rhs_element_index The index of the second element
-     * @return Indicates the relationship between the elements in the `lhs` and `rhs` columns, along
-     * with the depth at which a null value was encountered.
-     */
-    template <typename Element,
-              CUDF_ENABLE_IF(has_nested_columns and std::is_same_v<Element, cudf::list_view>)>
-    __device__ cuda::std::pair<weak_ordering, int> operator()(size_type lhs_element_index,
-                                                              size_type rhs_element_index)
-    {
-      // only order top-NULLs according to null_order
-      auto const is_l_row_null = _lhs.is_null(lhs_element_index);
-      auto const is_r_row_null = _rhs.is_null(rhs_element_index);
-      if (is_l_row_null || is_r_row_null) {
-        return cuda::std::pair(null_compare(is_l_row_null, is_r_row_null, _null_precedence),
-                               _depth);
-      }
-
-      // These are all the values from the Dremel encoding.
-      auto const l_max_def_level = _l_dremel_device_view->max_def_level;
-      auto const r_max_def_level = _r_dremel_device_view->max_def_level;
-      auto const l_def_levels    = _l_dremel_device_view->def_levels;
-      auto const r_def_levels    = _r_dremel_device_view->def_levels;
-      auto const l_rep_levels    = _l_dremel_device_view->rep_levels;
-      auto const r_rep_levels    = _r_dremel_device_view->rep_levels;
-
-      // Traverse the nested list hierarchy to get a column device view
-      // pointing to the underlying child data.
-      column_device_view lcol = _lhs.slice(lhs_element_index, 1);
-      column_device_view rcol = _rhs.slice(rhs_element_index, 1);
-
-      while (lcol.type().id() == type_id::LIST) {
-        lcol = detail::lists_column_device_view(lcol).get_sliced_child();
-        rcol = detail::lists_column_device_view(rcol).get_sliced_child();
-      }
-
-      // These start and end values indicate the start and end points of all
-      // the elements of the lists in the current list element
-      // (`[lhs|rhs]_element_index`) that we are comparing.
-      auto const l_offsets = _l_dremel_device_view->offsets;
-      auto const r_offsets = _r_dremel_device_view->offsets;
-      auto l_start         = l_offsets[lhs_element_index];
-      auto l_end           = l_offsets[lhs_element_index + 1];
-      auto r_start         = r_offsets[rhs_element_index];
-      auto r_end           = r_offsets[rhs_element_index + 1];
-
-      // This comparator will be used to compare leaf (non-nested) data types.
-      auto comparator =
-        element_comparator{_check_nulls, lcol, rcol, _null_precedence, _depth, _comparator};
-
-      // Loop over each element in the encoding. Note that this includes nulls
-      // and empty lists, so not every index corresponds to an actual element
-      // in the child column. The element_index is used to keep track of the current
-      // child element that we're actually comparing.
-      for (int l_dremel_index = l_start, r_dremel_index = r_start, element_index = 0;
-           l_dremel_index < l_end and r_dremel_index < r_end;
-           ++l_dremel_index, ++r_dremel_index) {
-        auto const l_rep_level = l_rep_levels[l_dremel_index];
-        auto const r_rep_level = r_rep_levels[r_dremel_index];
-
-        // early exit for smaller sub-list
-        if (l_rep_level != r_rep_level) {
-          // the lower repetition level is a smaller sub-list
-          return l_rep_level < r_rep_level ? cuda::std::pair(weak_ordering::LESS, _depth)
-                                           : cuda::std::pair(weak_ordering::GREATER, _depth);
-        }
-
-        // only compare if left and right are at same nesting level
-        auto const l_def_level = l_def_levels[l_dremel_index];
-        auto const r_def_level = r_def_levels[r_dremel_index];
-
-        // either left or right are empty or NULLs of arbitrary nesting
-        if (l_def_level < l_max_def_level || r_def_level < r_max_def_level) {
-          // in the fully unraveled version of the list column, only the
-          // most nested NULLs and leafs are present
-          // In this rare condition that we get to the most nested NULL, we increment
-          // element_index because either both rows have a deeply nested NULL at the
-          // same position, and we'll "continue" in our iteration, or we will early
-          // exit if only one of the rows has a deeply nested NULL
-          if ((lcol.nullable() and l_def_levels[l_dremel_index] == l_max_def_level - 1) or
-              (rcol.nullable() and r_def_levels[r_dremel_index] == r_max_def_level - 1)) {
-            ++element_index;
-          }
-          if (l_def_level == r_def_level) { continue; }
-          // We require [] < [NULL] < [leaf] for nested nulls.
-          // The null_precedence only affects top level nulls.
-          return l_def_level < r_def_level ? cuda::std::pair(weak_ordering::LESS, _depth)
-                                           : cuda::std::pair(weak_ordering::GREATER, _depth);
-        }
-
-        // finally, compare leaf to leaf
-        weak_ordering state{weak_ordering::EQUIVALENT};
-        int last_null_depth                    = _depth;
-        cuda::std::tie(state, last_null_depth) = cudf::type_dispatcher<dispatch_void_if_nested>(
-          lcol.type(), comparator, element_index, element_index);
-        if (state != weak_ordering::EQUIVALENT) { return cuda::std::pair(state, _depth); }
-        ++element_index;
-      }
-
-      // If we have reached this stage, we know that definition levels,
-      // repetition levels, and actual elements are identical in both list
-      // columns up to the `min(l_end - l_start, r_end - r_start)` element of
-      // the Dremel encoding. However, two lists can only compare equivalent if
-      // they are of the same length. Otherwise, the shorter of the two is less
-      // than the longer. This final check determines the appropriate resulting
-      // ordering by checking how many total elements each list is composed of.
-      return cuda::std::pair(detail::compare_elements(l_end - l_start, r_end - r_start), _depth);
-    }
-
-   private:
-    column_device_view const _lhs;
-    column_device_view const _rhs;
-    Nullate const _check_nulls;
-    null_order const _null_precedence;
-    int const _depth;
-    optional_dremel_view _l_dremel_device_view;
-    optional_dremel_view _r_dremel_device_view;
-    PhysicalElementComparator const _comparator;
-  };
-
- public:
-  /**
-   * @brief Checks whether the row at `lhs_index` in the `lhs` table compares
-   * lexicographically less, greater, or equivalent to the row at `rhs_index` in the `rhs` table.
-   *
-   * @param lhs_index The index of the row in the `lhs` table to examine
-   * @param rhs_index The index of the row in the `rhs` table to examine
-   * @return weak ordering comparison of the row in the `lhs` table relative to the row in the `rhs`
-   * table
-   */
-  __device__ constexpr weak_ordering operator()(size_type const lhs_index,
-                                                size_type const rhs_index) const noexcept
-  {
-    int last_null_depth = cuda::std::numeric_limits<int>::max();
-    size_type list_column_index{-1};
-    for (size_type i = 0; i < _lhs.num_columns(); ++i) {
-      if (_lhs.column(i).type().id() == type_id::LIST) { ++list_column_index; }
-
-      int const depth = _depth.has_value() ? (*_depth)[i] : 0;
-      if (depth > last_null_depth) { continue; }
-
-      bool const ascending =
-        _column_order.has_value() ? (*_column_order)[i] == order::ASCENDING : true;
-
-      null_order const null_precedence =
-        _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE;
-
-      // TODO: At what point do we verify that the columns of lhs and rhs are
-      // all of the same types? I assume that it's already happened before
-      // here, otherwise the current code would be failing.
-      auto const [l_dremel_i, r_dremel_i] =
-        _lhs.column(i).type().id() == type_id::LIST
-          ? cuda::std::make_tuple(optional_dremel_view(_l_dremel[list_column_index]),
-                                  optional_dremel_view(_r_dremel[list_column_index]))
-          : cuda::std::make_tuple(optional_dremel_view{}, optional_dremel_view{});
-
-      auto element_comp = element_comparator{_check_nulls,
-                                             _lhs.column(i),
-                                             _rhs.column(i),
-                                             null_precedence,
-                                             depth,
-                                             _comparator,
-                                             l_dremel_i,
-                                             r_dremel_i};
-
-      weak_ordering state;
-      cuda::std::tie(state, last_null_depth) =
-        cudf::type_dispatcher(_lhs.column(i).type(), element_comp, lhs_index, rhs_index);
-
-      if (state == weak_ordering::EQUIVALENT) { continue; }
-
-      return ascending
-               ? state
-               : (state == weak_ordering::GREATER ? weak_ordering::LESS : weak_ordering::GREATER);
-    }
-    return weak_ordering::EQUIVALENT;
-  }
-
- private:
-  table_device_view const _lhs;
-  table_device_view const _rhs;
-  device_span<detail::dremel_device_view const> const _l_dremel;
-  device_span<detail::dremel_device_view const> const _r_dremel;
-  Nullate const _check_nulls;
-  cuda::std::optional<device_span<int const>> const _depth;
-  cuda::std::optional<device_span<order const>> const _column_order;
-  cuda::std::optional<device_span<null_order const>> const _null_precedence;
-  PhysicalElementComparator const _comparator;
-};  // class device_row_comparator
-
-/**
- * @brief Wraps and interprets the result of templated Comparator that returns a weak_ordering.
- * Returns true if the weak_ordering matches any of the templated values.
- *
- * Note that this should never be used with only `weak_ordering::EQUIVALENT`.
- * An equality comparator should be used instead for optimal performance.
- *
- * @tparam Comparator generic comparator that returns a weak_ordering.
- * @tparam values weak_ordering parameter pack of orderings to interpret as true
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <typename Comparator, weak_ordering... values>
-struct weak_ordering_comparator_impl {
-  static_assert(not((weak_ordering::EQUIVALENT == values) && ...),
-                "weak_ordering_comparator should not be used for pure equality comparisons. The "
-                "`row_equality_comparator` should be used instead");
-
-  template <typename LhsType, typename RhsType>
-  __device__ constexpr bool operator()(LhsType const lhs_index,
-                                       RhsType const rhs_index) const noexcept
-  {
-    weak_ordering const result = comparator(lhs_index, rhs_index);
-    return ((result == values) || ...);
-  }
-  Comparator const comparator;
-};
-
-/**
- * @brief Wraps and interprets the result of device_row_comparator, true if the result is
- * weak_ordering::LESS meaning one row is lexicographically *less* than another row.
- *
- * @tparam Comparator generic comparator that returns a weak_ordering
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <typename Comparator>
-struct less_comparator : weak_ordering_comparator_impl<Comparator, weak_ordering::LESS> {
-  /**
-   * @brief Constructs a less_comparator
-   *
-   * @param comparator The comparator to wrap
-   */
-  less_comparator(Comparator const& comparator)
-    : weak_ordering_comparator_impl<Comparator, weak_ordering::LESS>{comparator}
-  {
-  }
-};
-
-/**
- * @brief Wraps and interprets the result of device_row_comparator, true if the result is
- * weak_ordering::LESS or weak_ordering::EQUIVALENT meaning one row is lexicographically *less* than
- * or *equivalent* to another row.
- *
- * @tparam Comparator generic comparator that returns a weak_ordering
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <typename Comparator>
-struct less_equivalent_comparator
-  : weak_ordering_comparator_impl<Comparator, weak_ordering::LESS, weak_ordering::EQUIVALENT> {
-  /**
-   * @brief Constructs a less_equivalent_comparator
-   *
-   * @param comparator The comparator to wrap
-   */
-  less_equivalent_comparator(Comparator const& comparator)
-    : weak_ordering_comparator_impl<Comparator, weak_ordering::LESS, weak_ordering::EQUIVALENT>{
-        comparator}
-  {
-  }
-};
-
-/**
- * @brief Preprocessed table for use with lexicographical comparison
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-struct preprocessed_table {
-  /// Type of table device view owner for the preprocessed table.
-  using table_device_view_owner =
-    std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
-
-  /**
-   * @brief Preprocess table for use with lexicographical comparison
-   *
-   * Sets up the table for use with lexicographical comparison. The resulting preprocessed table can
-   * be passed to the constructor of `lexicographic::self_comparator` or
-   * `lexicographic::two_table_comparator` to avoid preprocessing again.
-   *
-   * Note that the output of this factory function should not be used in `two_table_comparator` if
-   * the input table contains lists-of-structs. In such cases, please use the overload
-   * `preprocessed_table::create(table_view const&, table_view const&,...)` to preprocess both input
-   * tables at the same time.
-   *
-   * @param table The table to preprocess
-   * @param column_order Optional, host array the same length as a row that indicates the desired
-   *        ascending/descending order of each column in a row. If empty, it is assumed all columns
-   *        are sorted in ascending order.
-   * @param null_precedence Optional, an array having the same length as the number of columns in
-   *        the input tables that indicates how null values compare to all other. If it is empty,
-   *        the order `null_order::BEFORE` will be used for all columns.
-   * @param stream The stream to launch kernels and h->d copies on while preprocessing
-   * @return A shared pointer to a preprocessed table
-   */
-  static std::shared_ptr<preprocessed_table> create(table_view const& table,
-                                                    host_span<order const> column_order,
-                                                    host_span<null_order const> null_precedence,
-                                                    rmm::cuda_stream_view stream);
-
-  /**
-   * @brief Preprocess tables for use with lexicographical comparison
-   *
-   * Sets up the tables for use with lexicographical comparison. The resulting preprocessed tables
-   * can be passed to the constructor of `lexicographic::self_comparator` or
-   * `lexicographic::two_table_comparator` to avoid preprocessing again.
-   *
-   * This factory function performs some extra operations to guarantee that its output can be used
-   * in `two_table_comparator` for all cases.
-   *
-   * @param lhs The lhs table to preprocess
-   * @param rhs The rhs table to preprocess
-   * @param column_order Optional, host array the same length as a row that indicates the desired
-   *        ascending/descending order of each column in a row. If empty, it is assumed all columns
-   *        are sorted in ascending order.
-   * @param null_precedence Optional, an array having the same length as the number of columns in
-   *        the input tables that indicates how null values compare to all other. If it is empty,
-   *        the order `null_order::BEFORE` will be used for all columns.
-   * @param stream The stream to launch kernels and h->d copies on while preprocessing
-   * @return A pair of shared pointers to the preprocessed tables
-   */
-  static std::pair<std::shared_ptr<preprocessed_table>, std::shared_ptr<preprocessed_table>> create(
-    table_view const& lhs,
-    table_view const& rhs,
-    host_span<order const> column_order,
-    host_span<null_order const> null_precedence,
-    rmm::cuda_stream_view stream);
-
- private:
-  friend class self_comparator;       ///< Allow self_comparator to access private members
-  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
-
-  /**
-   * @brief Create the output preprocessed table from intermediate preprocessing results
-   *
-   * @param preprocessed_input The table resulted from preprocessing
-   * @param verticalized_col_depths The depths of each column resulting from decomposing struct
-   *        columns in the original input table
-   * @param transformed_columns Store the intermediate columns generated from transforming
-   *        nested children columns into integers columns using `cudf::rank()`
-   * @param column_order Optional, host array the same length as a row that indicates the desired
-   *        ascending/descending order of each column in a row. If empty, it is assumed all columns
-   *        are sorted in ascending order.
-   * @param null_precedence Optional, an array having the same length as the number of columns in
-   *        the input tables that indicates how null values compare to all other. If it is empty,
-   *        the order `null_order::BEFORE` will be used for all columns.
-   * @param has_ranked_children Flag indicating if the input table was preprocessed to transform
-   *        any nested child column into an integer column using `cudf::rank`
-   * @param stream The stream to launch kernels and h->d copies on while preprocessing
-   * @return A shared pointer to a preprocessed table
-   */
-  static std::shared_ptr<preprocessed_table> create(
-    table_view const& preprocessed_input,
-    std::vector<int>&& verticalized_col_depths,
-    std::vector<std::unique_ptr<column>>&& transformed_columns,
-    host_span<order const> column_order,
-    host_span<null_order const> null_precedence,
-    bool has_ranked_children,
-    rmm::cuda_stream_view stream);
-
-  /**
-   * @brief Construct a preprocessed table for use with lexicographical comparison
-   *
-   * Sets up the table for use with lexicographical comparison. The resulting preprocessed table can
-   * be passed to the constructor of `lexicographic::self_comparator` to avoid preprocessing again.
-   *
-   * @param table The table to preprocess
-   * @param column_order Optional, device array the same length as a row that indicates the desired
-   *        ascending/descending order of each column in a row. If empty, it is assumed all columns
-   *        are sorted in ascending order.
-   * @param null_precedence Optional, device array the same length as a row and indicates how null
-   *        values compare to all other for every column. If it is nullptr, then null precedence
-   *        would be `null_order::BEFORE` for all columns.
-   * @param depths The depths of each column resulting from decomposing struct columns.
-   * @param dremel_data The dremel data for each list column. The length of this object is the
-   *        number of list columns in the table.
-   * @param dremel_device_views Device views into the dremel_data structs contained in the
-   *        `dremel_data` parameter. For columns that are not list columns, this uvector will should
-   *        contain an empty `dremel_device_view`. As such, this uvector has as many elements as
-   *        there are columns in the table (unlike the `dremel_data` parameter, which is only as
-   *        long as the number of list columns).
-   * @param transformed_columns Store the intermediate columns generated from transforming
-   *        nested children columns into integers columns using `cudf::rank()`
-   * @param has_ranked_children Flag indicating if the input table was preprocessed to transform
-   *        any lists-of-structs column having floating-point children using `cudf::rank`
-   */
-  preprocessed_table(table_device_view_owner&& table,
-                     rmm::device_uvector<order>&& column_order,
-                     rmm::device_uvector<null_order>&& null_precedence,
-                     rmm::device_uvector<size_type>&& depths,
-                     std::vector<detail::dremel_data>&& dremel_data,
-                     rmm::device_uvector<detail::dremel_device_view>&& dremel_device_views,
-                     std::vector<std::unique_ptr<column>>&& transformed_columns,
-                     bool has_ranked_children);
-
-  preprocessed_table(table_device_view_owner&& table,
-                     rmm::device_uvector<order>&& column_order,
-                     rmm::device_uvector<null_order>&& null_precedence,
-                     rmm::device_uvector<size_type>&& depths,
-                     std::vector<std::unique_ptr<column>>&& transformed_columns,
-                     bool has_ranked_children);
-
-  /**
-   * @brief Implicit conversion operator to a `table_device_view` of the preprocessed table.
-   *
-   * @return table_device_view
-   */
-  operator table_device_view() { return *_t; }
-
-  /**
-   * @brief Get a device array containing the desired order of each column in the preprocessed table
-   *
-   * @return Device array containing respective column orders. If no explicit column orders were
-   * specified during the creation of this object then this will be `nullopt`.
-   */
-  [[nodiscard]] cuda::std::optional<device_span<order const>> column_order() const
-  {
-    return _column_order.size() ? cuda::std::optional<device_span<order const>>(_column_order)
-                                : cuda::std::nullopt;
-  }
-
-  /**
-   * @brief Get a device array containing the desired null precedence of each column in the
-   * preprocessed table
-   *
-   * @return Device array containing respective column null precedence. If no explicit column null
-   * precedences were specified during the creation of this object then this will be `nullopt`.
-   */
-  [[nodiscard]] cuda::std::optional<device_span<null_order const>> null_precedence() const
-  {
-    return _null_precedence.size()
-             ? cuda::std::optional<device_span<null_order const>>(_null_precedence)
-             : cuda::std::nullopt;
-  }
-
-  /**
-   * @brief Get a device array containing the depth of each column in the preprocessed table
-   *
-   * @see struct_linearize()
-   *
-   * @return std::optional<device_span<int const>> Device array containing respective column depths.
-   * If there are no nested columns in the table then this will be `nullopt`.
-   */
-  [[nodiscard]] cuda::std::optional<device_span<int const>> depths() const
-  {
-    return _depths.size() ? cuda::std::optional<device_span<int const>>(_depths)
-                          : cuda::std::nullopt;
-  }
-
-  [[nodiscard]] device_span<detail::dremel_device_view const> dremel_device_views() const
-  {
-    if (_dremel_device_views.has_value()) {
-      return device_span<detail::dremel_device_view const>(*_dremel_device_views);
-    } else {
-      return {};
-    }
-  }
-
-  template <typename PhysicalElementComparator>
-  void check_physical_element_comparator()
-  {
-    if constexpr (!std::is_same_v<PhysicalElementComparator, sorting_physical_element_comparator>) {
-      CUDF_EXPECTS(!_has_ranked_children,
-                   "The input table has nested type children and they were transformed using a "
-                   "different type of physical element comparator.");
-    }
-  }
-
- private:
-  table_device_view_owner const _t;
-  rmm::device_uvector<order> const _column_order;
-  rmm::device_uvector<null_order> const _null_precedence;
-  rmm::device_uvector<size_type> const _depths;
-
-  // Dremel encoding of list columns used for the comparison algorithm
-  cuda::std::optional<std::vector<detail::dremel_data>> _dremel_data;
-  cuda::std::optional<rmm::device_uvector<detail::dremel_device_view>> _dremel_device_views;
-
-  // Intermediate columns generated from transforming nested children columns into
-  // integers columns using `cudf::rank()`, need to be kept alive.
-  std::vector<std::unique_ptr<column>> _transformed_columns;
-
-  // Flag to record if the input table was preprocessed to transform any nested children column(s)
-  // into integer column(s) using `cudf::rank`.
-  bool const _has_ranked_children;
-};
-
-/**
- * @brief An owning object that can be used to lexicographically compare two rows of the same table
- *
- * This class can take a table_view and preprocess certain columns to allow for lexicographical
- * comparison. The preprocessed table and temporary data required for the comparison are created and
- * owned by this class.
- *
- * Alternatively, `self_comparator` can be constructed from an existing
- * `shared_ptr<preprocessed_table>` when sharing the same table among multiple comparators.
- *
- * This class can then provide a functor object that can used on the device.
- * The object of this class must outlive the usage of the device functor.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-class self_comparator {
- public:
-  /**
-   * @brief Construct an owning object for performing a lexicographic comparison between two rows of
-   * the same table.
-   *
-   * @param t The table to compare
-   * @param column_order Optional, host array the same length as a row that indicates the desired
-   *        ascending/descending order of each column in a row. If empty, it is assumed all columns
-   *        are sorted in ascending order.
-   * @param null_precedence Optional, device array the same length as a row and indicates how null
-   *        values compare to all other for every column. If empty, then null precedence would be
-   *        `null_order::BEFORE` for all columns.
-   * @param stream The stream to construct this object on. Not the stream that will be used for
-   *        comparisons using this object.
-   */
-  self_comparator(table_view const& t,
-                  host_span<order const> column_order         = {},
-                  host_span<null_order const> null_precedence = {},
-                  rmm::cuda_stream_view stream                = cudf::get_default_stream())
-    : d_t{preprocessed_table::create(t, column_order, null_precedence, stream)}
-  {
-  }
-
-  /**
-   * @brief Construct an owning object for performing a lexicographic comparison between two rows of
-   * the same preprocessed table.
-   *
-   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
-   * multiple comparators.
-   *
-   * @param t A table preprocessed for lexicographic comparison
-   */
-  self_comparator(std::shared_ptr<preprocessed_table> t) : d_t{std::move(t)} {}
-
-  /**
-   * @brief Return the binary operator for comparing rows in the table.
-   *
-   * Returns a binary callable, `F`, with signature `bool F(size_type, size_type)`.
-   *
-   * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`.
-   *
-   * @note The operator overloads in sub-class `element_comparator` are templated via the
-   *       `type_dispatcher` to help select an overload instance for each column in a table.
-   *       So, `cudf::is_nested<Element>` will return `true` if the table has nested-type columns,
-   *       but it will be a runtime error if template parameter `has_nested_columns != true`.
-   *
-   * @tparam has_nested_columns compile-time optimization for primitive types.
-   *         This template parameter is to be used by the developer by querying
-   *         `cudf::has_nested_columns(input)`. `true` compiles operator
-   *         overloads for nested types, while `false` only compiles operator
-   *         overloads for primitive types.
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
-   * @tparam PhysicalElementComparator A relational comparator functor that compares individual
-   *         values rather than logical elements, defaults to `NaN` aware relational comparator
-   *         that evaluates `NaN` as greater than all other values.
-   * @throw cudf::logic_error if the input table was preprocessed to transform any nested children
-   *        columns into integer columns but `PhysicalElementComparator` is not
-   *        `sorting_physical_element_comparator`.
-   * @param nullate Indicates if any input column contains nulls.
-   * @param comparator Physical element relational comparison functor.
-   * @return A binary callable object.
-   */
-  template <bool has_nested_columns,
-            typename Nullate,
-            typename PhysicalElementComparator = sorting_physical_element_comparator>
-  auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const
-  {
-    d_t->check_physical_element_comparator<PhysicalElementComparator>();
-
-    return less_comparator{
-      device_row_comparator<has_nested_columns, Nullate, PhysicalElementComparator>{
-        nullate,
-        *d_t,
-        *d_t,
-        d_t->dremel_device_views(),
-        d_t->dremel_device_views(),
-        d_t->depths(),
-        d_t->column_order(),
-        d_t->null_precedence(),
-        comparator}};
-  }
-
-  /// @copydoc less()
-  template <bool has_nested_columns,
-            typename Nullate,
-            typename PhysicalElementComparator = sorting_physical_element_comparator>
-  auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const
-  {
-    d_t->check_physical_element_comparator<PhysicalElementComparator>();
-
-    return less_equivalent_comparator{
-      device_row_comparator<has_nested_columns, Nullate, PhysicalElementComparator>{
-        nullate,
-        *d_t,
-        *d_t,
-        d_t->dremel_device_views(),
-        d_t->dremel_device_views(),
-        d_t->depths(),
-        d_t->column_order(),
-        d_t->null_precedence(),
-        comparator}};
-  }
-
- private:
-  std::shared_ptr<preprocessed_table> d_t;
-};
-
-// @cond
-template <typename Comparator>
-struct strong_index_comparator_adapter {
-  strong_index_comparator_adapter(Comparator const& comparator) : comparator{comparator} {}
-
-  __device__ constexpr weak_ordering operator()(lhs_index_type const lhs_index,
-                                                rhs_index_type const rhs_index) const noexcept
-  {
-    return comparator(static_cast<cudf::size_type>(lhs_index),
-                      static_cast<cudf::size_type>(rhs_index));
-  }
-
-  __device__ constexpr weak_ordering operator()(rhs_index_type const rhs_index,
-                                                lhs_index_type const lhs_index) const noexcept
-  {
-    auto const left_right_ordering =
-      comparator(static_cast<cudf::size_type>(lhs_index), static_cast<cudf::size_type>(rhs_index));
-
-    // Invert less/greater values to reflect right to left ordering
-    if (left_right_ordering == weak_ordering::LESS) {
-      return weak_ordering::GREATER;
-    } else if (left_right_ordering == weak_ordering::GREATER) {
-      return weak_ordering::LESS;
-    }
-    return weak_ordering::EQUIVALENT;
-  }
-
-  Comparator const comparator;
-};
-// @endcond
-
-/**
- * @brief An owning object that can be used to lexicographically compare rows of two different
- * tables
- *
- * This class takes two table_views and preprocesses certain columns to allow for lexicographical
- * comparison. The preprocessed table and temporary data required for the comparison are created and
- * owned by this class.
- *
- * Alternatively, `two_table_comparator` can be constructed from two existing
- * `shared_ptr<preprocessed_table>`s when sharing the same tables among multiple comparators.
- *
- * This class can then provide a functor object that can used on the device.
- * The object of this class must outlive the usage of the device functor.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-class two_table_comparator {
- public:
-  /**
-   * @brief Construct an owning object for performing a lexicographic comparison between rows of
-   * two different tables.
-   *
-   * The left and right table are expected to have the same number of columns
-   * and data types for each column.
-   *
-   * @param left The left table to compare
-   * @param right The right table to compare
-   * @param column_order Optional, host array the same length as a row that indicates the desired
-   *        ascending/descending order of each column in a row. If empty, it is assumed all columns
-   *        are sorted in ascending order.
-   * @param null_precedence Optional, device array the same length as a row and indicates how null
-   *        values compare to all other for every column. If empty, then null precedence would be
-   *        `null_order::BEFORE` for all columns.
-   * @param stream The stream to construct this object on. Not the stream that will be used for
-   *        comparisons using this object.
-   */
-  two_table_comparator(table_view const& left,
-                       table_view const& right,
-                       host_span<order const> column_order         = {},
-                       host_span<null_order const> null_precedence = {},
-                       rmm::cuda_stream_view stream                = cudf::get_default_stream());
-
-  /**
-   * @brief Construct an owning object for performing a lexicographic comparison between two rows of
-   * the same preprocessed table.
-   *
-   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
-   * multiple comparators.
-   *
-   * The preprocessed_table(s) should have been pre-generated together using the factory function
-   * `preprocessed_table::create(table_view const&, table_view const&)`. Otherwise, the comparison
-   * results between two tables may be incorrect.
-   *
-   * @param left A table preprocessed for lexicographic comparison
-   * @param right A table preprocessed for lexicographic comparison
-   */
-  two_table_comparator(std::shared_ptr<preprocessed_table> left,
-                       std::shared_ptr<preprocessed_table> right)
-    : d_left_table{std::move(left)}, d_right_table{std::move(right)}
-  {
-  }
-
-  /**
-   * @brief Return the binary operator for comparing rows in the table.
-   *
-   * Returns a binary callable, `F`, with signatures
-   * `bool F(lhs_index_type, rhs_index_type)` and
-   * `bool F(rhs_index_type, lhs_index_type)`.
-   *
-   * `F(lhs_index_type i, rhs_index_type j)` returns true if and only if row
-   * `i` of the left table compares lexicographically less than row `j` of the
-   * right table.
-   *
-   * Similarly, `F(rhs_index_type i, lhs_index_type j)` returns true if and
-   * only if row `i` of the right table compares lexicographically less than row
-   * `j` of the left table.
-   *
-   * @note The operator overloads in sub-class `element_comparator` are templated via the
-   *       `type_dispatcher` to help select an overload instance for each column in a table.
-   *       So, `cudf::is_nested<Element>` will return `true` if the table has nested-type columns,
-   *       but it will be a runtime error if template parameter `has_nested_columns != true`.
-   *
-   * @tparam has_nested_columns compile-time optimization for primitive types.
-   *         This template parameter is to be used by the developer by querying
-   *         `cudf::has_nested_columns(input)`. `true` compiles operator
-   *         overloads for nested types, while `false` only compiles operator
-   *         overloads for primitive types.
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
-   * @tparam PhysicalElementComparator A relational comparator functor that compares individual
-   *         values rather than logical elements, defaults to `NaN` aware relational comparator
-   *         that evaluates `NaN` as greater than all other values.
-   * @throw cudf::logic_error if the input tables were preprocessed to transform any nested children
-   *        columns into integer columns but `PhysicalElementComparator` is not
-   *        `sorting_physical_element_comparator`.
-   * @param nullate Indicates if any input column contains nulls.
-   * @param comparator Physical element relational comparison functor.
-   * @return A binary callable object.
-   */
-  template <bool has_nested_columns,
-            typename Nullate,
-            typename PhysicalElementComparator = sorting_physical_element_comparator>
-  auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const
-  {
-    d_left_table->check_physical_element_comparator<PhysicalElementComparator>();
-    d_right_table->check_physical_element_comparator<PhysicalElementComparator>();
-
-    return less_comparator{strong_index_comparator_adapter{
-      device_row_comparator<has_nested_columns, Nullate, PhysicalElementComparator>{
-        nullate,
-        *d_left_table,
-        *d_right_table,
-        d_left_table->dremel_device_views(),
-        d_right_table->dremel_device_views(),
-        d_left_table->depths(),
-        d_left_table->column_order(),
-        d_left_table->null_precedence(),
-        comparator}}};
-  }
-
-  /// @copydoc less()
-  template <bool has_nested_columns,
-            typename Nullate,
-            typename PhysicalElementComparator = sorting_physical_element_comparator>
-  auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const
-  {
-    d_left_table->check_physical_element_comparator<PhysicalElementComparator>();
-    d_right_table->check_physical_element_comparator<PhysicalElementComparator>();
-
-    return less_equivalent_comparator{strong_index_comparator_adapter{
-      device_row_comparator<has_nested_columns, Nullate, PhysicalElementComparator>{
-        nullate,
-        *d_left_table,
-        *d_right_table,
-        d_left_table->dremel_device_views(),
-        d_right_table->dremel_device_views(),
-        d_left_table->depths(),
-        d_left_table->column_order(),
-        d_left_table->null_precedence(),
-        comparator}}};
-  }
-
- private:
-  std::shared_ptr<preprocessed_table> d_left_table;
-  std::shared_ptr<preprocessed_table> d_right_table;
-};
-
-}  // namespace lexicographic
-
-namespace hash {
-class row_hasher;
-}  // namespace hash
-
-namespace equality {
-
-/**
- * @brief Equality comparator functor that compares physical values rather than logical
- * elements like lists, strings, or structs. It evaluates `NaN` not equal to all other values for
- * IEEE-754 compliance.
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-struct physical_equality_comparator {
-  /**
-   * @brief Operator for equality comparisons.
-   *
-   * Note that `NaN != NaN`, following IEEE-754.
-   *
-   * @param lhs First element
-   * @param rhs Second element
-   * @return `true` if `lhs == rhs` else `false`
-   */
-  template <typename Element>
-  __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept
-  {
-    return lhs == rhs;
-  }
-};
-
-/**
- * @brief Equality comparator functor that compares physical values rather than logical
- * elements like lists, strings, or structs. It evaluates `NaN` as equal to other `NaN`s.
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-struct nan_equal_physical_equality_comparator {
-  /**
-   * @brief Operator for equality comparison of non-floating point values.
-   *
-   * @param lhs First element
-   * @param rhs Second element
-   * @return `true` if `lhs == rhs` else `false`
-   */
-  template <typename Element, CUDF_ENABLE_IF(not std::is_floating_point_v<Element>)>
-  __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept
-  {
-    return lhs == rhs;
-  }
-
-  /**
-   * @brief Operator for equality comparison of floating point values.
-   *
-   * Note that `NaN == NaN`.
-   *
-   * @param lhs First element
-   * @param rhs Second element
-   * @return `true` if `lhs` == `rhs` else `false`
-   */
-  template <typename Element, CUDF_ENABLE_IF(std::is_floating_point_v<Element>)>
-  __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept
-  {
-    return isnan(lhs) and isnan(rhs) ? true : lhs == rhs;
-  }
-};
-
-/**
- * @brief Computes the equality comparison between 2 rows.
- *
- * Equality is determined by comparing rows element by element. The first mismatching element
- * returns false, representing unequal rows. If the rows are compared without mismatched elements,
- * the rows are equal.
- *
- * @note The operator overloads in sub-class `element_comparator` are templated via the
- *        `type_dispatcher` to help select an overload instance for each column in a table.
- *        So, `cudf::is_nested<Element>` will return `true` if the table has nested-type columns,
- *        but it will be a runtime error if template parameter `has_nested_columns != true`.
- *
- * @tparam has_nested_columns compile-time optimization for primitive types.
- *         This template parameter is to be used by the developer by querying
- *         `cudf::has_nested_columns(input)`. `true` compiles operator
- *         overloads for nested types, while `false` only compiles operator
- *         overloads for primitive types.
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual values
- * rather than logical elements, defaults to a comparator for which `NaN == NaN`.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <bool has_nested_columns,
-          typename Nullate,
-          typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator>
-class device_row_comparator {
-  friend class self_comparator;       ///< Allow self_comparator to access private members
-  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
-
- public:
-  /**
-   * @brief Checks whether the row at `lhs_index` in the `lhs` table is equal to the row at
-   * `rhs_index` in the `rhs` table.
-   *
-   * @param lhs_index The index of the row in the `lhs` table to examine
-   * @param rhs_index The index of the row in the `rhs` table to examine
-   * @return `true` if row from the `lhs` table is equal to the row in the `rhs` table
-   */
-  __device__ constexpr bool operator()(size_type const lhs_index,
-                                       size_type const rhs_index) const noexcept
-  {
-    auto equal_elements = [lhs_index, rhs_index, this](column_device_view l, column_device_view r) {
-      return cudf::type_dispatcher(
-        l.type(),
-        element_comparator{check_nulls, l, r, nulls_are_equal, comparator},
-        lhs_index,
-        rhs_index);
-    };
-
-    return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements);
-  }
-
- private:
-  /**
-   * @brief Construct a function object for performing equality comparison between the rows of two
-   * tables.
-   *
-   * @param check_nulls Indicates if any input column contains nulls.
-   * @param lhs The first table
-   * @param rhs The second table (may be the same table as `lhs`)
-   * @param nulls_are_equal Indicates if two null elements are treated as equivalent
-   * @param comparator Physical element equality comparison functor.
-   */
-  device_row_comparator(Nullate check_nulls,
-                        table_device_view lhs,
-                        table_device_view rhs,
-                        null_equality nulls_are_equal         = null_equality::EQUAL,
-                        PhysicalEqualityComparator comparator = {}) noexcept
-    : lhs{lhs},
-      rhs{rhs},
-      check_nulls{check_nulls},
-      nulls_are_equal{nulls_are_equal},
-      comparator{comparator}
-  {
-  }
-
-  /**
-   * @brief Performs an equality comparison between two elements in two columns.
-   */
-  class element_comparator {
-   public:
-    /**
-     * @brief Construct type-dispatched function object for comparing equality
-     * between two elements.
-     *
-     * @note `lhs` and `rhs` may be the same.
-     *
-     * @param check_nulls Indicates if either input column contains nulls.
-     * @param lhs The column containing the first element
-     * @param rhs The column containing the second element (may be the same as lhs)
-     * @param nulls_are_equal Indicates if two null elements are treated as equivalent
-     * @param comparator Physical element equality comparison functor.
-     */
-    __device__ element_comparator(Nullate check_nulls,
-                                  column_device_view lhs,
-                                  column_device_view rhs,
-                                  null_equality nulls_are_equal         = null_equality::EQUAL,
-                                  PhysicalEqualityComparator comparator = {}) noexcept
-      : lhs{lhs},
-        rhs{rhs},
-        check_nulls{check_nulls},
-        nulls_are_equal{nulls_are_equal},
-        comparator{comparator}
-    {
-    }
-
-    /**
-     * @brief Compares the specified elements for equality.
-     *
-     * @param lhs_element_index The index of the first element
-     * @param rhs_element_index The index of the second element
-     * @return True if lhs and rhs are equal or if both lhs and rhs are null and nulls are
-     * considered equal (`nulls_are_equal` == `null_equality::EQUAL`)
-     */
-    template <typename Element, CUDF_ENABLE_IF(cudf::is_equality_comparable<Element, Element>())>
-    __device__ bool operator()(size_type const lhs_element_index,
-                               size_type const rhs_element_index) const noexcept
-    {
-      if (check_nulls) {
-        bool const lhs_is_null{lhs.is_null(lhs_element_index)};
-        bool const rhs_is_null{rhs.is_null(rhs_element_index)};
-        if (lhs_is_null and rhs_is_null) {
-          return nulls_are_equal == null_equality::EQUAL;
-        } else if (lhs_is_null != rhs_is_null) {
-          return false;
-        }
-      }
-
-      return comparator(lhs.element<Element>(lhs_element_index),
-                        rhs.element<Element>(rhs_element_index));
-    }
-
-    template <typename Element,
-              CUDF_ENABLE_IF(not cudf::is_equality_comparable<Element, Element>() and
-                             (not has_nested_columns or not cudf::is_nested<Element>())),
-              typename... Args>
-    __device__ bool operator()(Args...)
-    {
-      CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
-    }
-
-    template <typename Element, CUDF_ENABLE_IF(has_nested_columns and cudf::is_nested<Element>())>
-    __device__ bool operator()(size_type const lhs_element_index,
-                               size_type const rhs_element_index) const noexcept
-    {
-      column_device_view lcol = lhs.slice(lhs_element_index, 1);
-      column_device_view rcol = rhs.slice(rhs_element_index, 1);
-      while (lcol.type().id() == type_id::STRUCT || lcol.type().id() == type_id::LIST) {
-        if (check_nulls) {
-          auto lvalid = detail::make_validity_iterator<true>(lcol);
-          auto rvalid = detail::make_validity_iterator<true>(rcol);
-          if (nulls_are_equal == null_equality::UNEQUAL) {
-            if (thrust::any_of(
-                  thrust::seq, lvalid, lvalid + lcol.size(), cuda::std::logical_not<bool>()) or
-                thrust::any_of(
-                  thrust::seq, rvalid, rvalid + rcol.size(), cuda::std::logical_not<bool>())) {
-              return false;
-            }
-          } else {
-            if (not thrust::equal(thrust::seq, lvalid, lvalid + lcol.size(), rvalid)) {
-              return false;
-            }
-          }
-        }
-        if (lcol.type().id() == type_id::STRUCT) {
-          if (lcol.num_child_columns() == 0) { return true; }
-          // Non-empty structs are assumed to be decomposed and contain only one child
-          lcol = detail::structs_column_device_view(lcol).get_sliced_child(0);
-          rcol = detail::structs_column_device_view(rcol).get_sliced_child(0);
-        } else if (lcol.type().id() == type_id::LIST) {
-          auto l_list_col = detail::lists_column_device_view(lcol);
-          auto r_list_col = detail::lists_column_device_view(rcol);
-
-          auto lsizes = make_list_size_iterator(l_list_col);
-          auto rsizes = make_list_size_iterator(r_list_col);
-          if (not thrust::equal(thrust::seq, lsizes, lsizes + lcol.size(), rsizes)) {
-            return false;
-          }
-
-          lcol = l_list_col.get_sliced_child();
-          rcol = r_list_col.get_sliced_child();
-          if (lcol.size() != rcol.size()) { return false; }
-        }
-      }
-
-      auto comp = column_comparator{
-        element_comparator{check_nulls, lcol, rcol, nulls_are_equal, comparator}, lcol.size()};
-      return type_dispatcher<dispatch_void_if_nested>(lcol.type(), comp);
-    }
-
-   private:
-    /**
-     * @brief Serially compare two columns for equality.
-     *
-     * When we want to get the equivalence of two columns by serially comparing all elements in
-     * one column with the corresponding elements in the other column, this saves us from type
-     * dispatching for each individual element in the range
-     */
-    struct column_comparator {
-      element_comparator const comp;
-      size_type const size;
-
-      /**
-       * @brief Serially compare two columns for equality.
-       *
-       * @return True if ALL elements compare equal, false otherwise
-       */
-      template <typename Element, CUDF_ENABLE_IF(cudf::is_equality_comparable<Element, Element>())>
-      __device__ bool operator()() const noexcept
-      {
-        return thrust::all_of(thrust::seq,
-                              thrust::make_counting_iterator(0),
-                              thrust::make_counting_iterator(0) + size,
-                              [this](auto i) { return comp.template operator()<Element>(i, i); });
-      }
-
-      template <typename Element,
-                CUDF_ENABLE_IF(not cudf::is_equality_comparable<Element, Element>()),
-                typename... Args>
-      __device__ bool operator()(Args...) const noexcept
-      {
-        CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
-      }
-    };
-
-    column_device_view const lhs;
-    column_device_view const rhs;
-    Nullate const check_nulls;
-    null_equality const nulls_are_equal;
-    PhysicalEqualityComparator const comparator;
-  };
-
-  table_device_view const lhs;
-  table_device_view const rhs;
-  Nullate const check_nulls;
-  null_equality const nulls_are_equal;
-  PhysicalEqualityComparator const comparator;
-};
-
-/**
- * @brief Preprocessed table for use with row equality comparison or row hashing
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-struct preprocessed_table {
-  /**
-   * @brief Factory to construct preprocessed_table for use with
-   * row equality comparison or row hashing
-   *
-   * Sets up the table for use with row equality comparison or row hashing. The resulting
-   * preprocessed table can be passed to the constructor of `equality::self_comparator` to
-   * avoid preprocessing again.
-   *
-   * @param table The table to preprocess
-   * @param stream The cuda stream to use while preprocessing.
-   * @return A preprocessed table as shared pointer
-   */
-  static std::shared_ptr<preprocessed_table> create(table_view const& table,
-                                                    rmm::cuda_stream_view stream);
-
- private:
-  friend class self_comparator;       ///< Allow self_comparator to access private members
-  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
-  friend class hash::row_hasher;      ///< Allow row_hasher to access private members
-  /// Allow primitive equality comparator to access private members
-  friend class ::cudf::row::primitive::row_equality_comparator;
-
-  template <template <typename> class Hash>
-  friend class ::cudf::row::primitive::row_hasher;
-
-  using table_device_view_owner =
-    std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
-
-  preprocessed_table(table_device_view_owner&& table,
-                     std::vector<rmm::device_buffer>&& null_buffers,
-                     std::vector<std::unique_ptr<column>>&& tmp_columns)
-    : _t(std::move(table)),
-      _null_buffers(std::move(null_buffers)),
-      _tmp_columns(std::move(tmp_columns))
-  {
-  }
-
-  /**
-   * @brief Implicit conversion operator to a `table_device_view` of the preprocessed table.
-   *
-   * @return table_device_view
-   */
-  operator table_device_view() { return *_t; }
-
-  table_device_view_owner _t;
-  std::vector<rmm::device_buffer> _null_buffers;
-  std::vector<std::unique_ptr<column>> _tmp_columns;
-};
-
-/**
- * @brief Comparator for performing equality comparisons between two rows of the same table.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-class self_comparator {
- public:
-  /**
-   * @brief Construct an owning object for performing equality comparisons between two rows of the
-   * same table.
-   *
-   * @param t The table to compare
-   * @param stream The stream to construct this object on. Not the stream that will be used for
-   * comparisons using this object.
-   */
-  self_comparator(table_view const& t, rmm::cuda_stream_view stream)
-    : d_t(preprocessed_table::create(t, stream))
-  {
-  }
-
-  /**
-   * @brief Construct an owning object for performing equality comparisons between two rows of the
-   * same table.
-   *
-   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
-   * multiple comparators.
-   *
-   * @param t A table preprocessed for equality comparison
-   */
-  self_comparator(std::shared_ptr<preprocessed_table> t) : d_t{std::move(t)} {}
-
-  /**
-   * @brief Get the comparison operator to use on the device
-   *
-   * Returns a binary callable, `F`, with signature `bool F(size_type, size_type)`.
-   *
-   * `F(i,j)` returns true if and only if row `i` compares equal to row `j`.
-   *
-   * @note The operator overloads in sub-class `element_comparator` are templated via the
-   *        `type_dispatcher` to help select an overload instance for each column in a table.
-   *        So, `cudf::is_nested<Element>` will return `true` if the table has nested-type columns,
-   *        but it will be a runtime error if template parameter `has_nested_columns != true`.
-   *
-   * @tparam has_nested_columns compile-time optimization for primitive types.
-   *         This template parameter is to be used by the developer by querying
-   *         `cudf::has_nested_columns(input)`. `true` compiles operator
-   *         overloads for nested types, while `false` only compiles operator
-   *         overloads for primitive types.
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
-   * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual
-   * values rather than logical elements, defaults to a comparator for which `NaN == NaN`.
-   * @param nullate Indicates if any input column contains nulls.
-   * @param nulls_are_equal Indicates if nulls are equal.
-   * @param comparator Physical element equality comparison functor.
-   * @return A binary callable object
-   */
-  template <bool has_nested_columns,
-            typename Nullate,
-            typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator>
-  auto equal_to(Nullate nullate                       = {},
-                null_equality nulls_are_equal         = null_equality::EQUAL,
-                PhysicalEqualityComparator comparator = {}) const noexcept
-  {
-    return device_row_comparator<has_nested_columns, Nullate, PhysicalEqualityComparator>{
-      nullate, *d_t, *d_t, nulls_are_equal, comparator};
-  }
-
- private:
-  std::shared_ptr<preprocessed_table> d_t;
-};
-
-// @cond
-template <typename Comparator>
-struct strong_index_comparator_adapter {
-  strong_index_comparator_adapter(Comparator const& comparator) : comparator{comparator} {}
-
-  __device__ constexpr bool operator()(lhs_index_type const lhs_index,
-                                       rhs_index_type const rhs_index) const noexcept
-  {
-    return comparator(static_cast<cudf::size_type>(lhs_index),
-                      static_cast<cudf::size_type>(rhs_index));
-  }
-
-  __device__ constexpr bool operator()(rhs_index_type const rhs_index,
-                                       lhs_index_type const lhs_index) const noexcept
-  {
-    return this->operator()(lhs_index, rhs_index);
-  }
-
-  Comparator const comparator;
-};
-// @endcond
-
-/**
- * @brief An owning object that can be used to equality compare rows of two different tables.
- *
- * This class takes two table_views and preprocesses certain columns to allow for equality
- * comparison. The preprocessed table and temporary data required for the comparison are created and
- * owned by this class.
- *
- * Alternatively, `two_table_comparator` can be constructed from two existing
- * `shared_ptr<preprocessed_table>`s when sharing the same tables among multiple comparators.
- *
- * This class can then provide a functor object that can used on the device.
- * The object of this class must outlive the usage of the device functor.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-class two_table_comparator {
- public:
-  /**
-   * @brief Construct an owning object for performing equality comparisons between two rows from two
-   * tables.
-   *
-   * The left and right table are expected to have the same number of columns and data types for
-   * each column.
-   *
-   * @param left The left table to compare.
-   * @param right The right table to compare.
-   * @param stream The stream to construct this object on. Not the stream that will be used for
-   * comparisons using this object.
-   */
-  two_table_comparator(table_view const& left,
-                       table_view const& right,
-                       rmm::cuda_stream_view stream);
-
-  /**
-   * @brief Construct an owning object for performing equality comparisons between two rows from two
-   * tables.
-   *
-   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
-   * multiple comparators.
-   *
-   * @param left The left table preprocessed for equality comparison.
-   * @param right The right table preprocessed for equality comparison.
-   */
-  two_table_comparator(std::shared_ptr<preprocessed_table> left,
-                       std::shared_ptr<preprocessed_table> right)
-    : d_left_table{std::move(left)}, d_right_table{std::move(right)}
-  {
-  }
-
-  /**
-   * @brief Return the binary operator for comparing rows in the table.
-   *
-   * Returns a binary callable, `F`, with signatures `bool F(lhs_index_type, rhs_index_type)` and
-   * `bool F(rhs_index_type, lhs_index_type)`.
-   *
-   * `F(lhs_index_type i, rhs_index_type j)` returns true if and only if row `i` of the left table
-   * compares equal to row `j` of the right table.
-   *
-   * Similarly, `F(rhs_index_type i, lhs_index_type j)` returns true if and only if row `i` of the
-   * right table compares equal to row `j` of the left table.
-   *
-   * @note The operator overloads in sub-class `element_comparator` are templated via the
-   *        `type_dispatcher` to help select an overload instance for each column in a table.
-   *        So, `cudf::is_nested<Element>` will return `true` if the table has nested-type columns,
-   *        but it will be a runtime error if template parameter `has_nested_columns != true`.
-   *
-   * @tparam has_nested_columns compile-time optimization for primitive types.
-   *         This template parameter is to be used by the developer by querying
-   *         `cudf::has_nested_columns(input)`. `true` compiles operator
-   *         overloads for nested types, while `false` only compiles operator
-   *         overloads for primitive types.
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
-   * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual
-   * values rather than logical elements, defaults to a `NaN == NaN` equality comparator.
-   * @param nullate Indicates if any input column contains nulls.
-   * @param nulls_are_equal Indicates if nulls are equal.
-   * @param comparator Physical element equality comparison functor.
-   * @return A binary callable object
-   */
-  template <bool has_nested_columns,
-            typename Nullate,
-            typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator>
-  auto equal_to(Nullate nullate                       = {},
-                null_equality nulls_are_equal         = null_equality::EQUAL,
-                PhysicalEqualityComparator comparator = {}) const noexcept
-  {
-    return strong_index_comparator_adapter{
-      device_row_comparator<has_nested_columns, Nullate, PhysicalEqualityComparator>(
-        nullate, *d_left_table, *d_right_table, nulls_are_equal, comparator)};
-  }
-
- private:
-  std::shared_ptr<preprocessed_table> d_left_table;
-  std::shared_ptr<preprocessed_table> d_right_table;
-};
-
-}  // namespace equality
-
-namespace hash {
-
-/**
- * @brief Computes the hash value of an element in the given column.
- *
- * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <template <typename> class hash_function, typename Nullate>
-class element_hasher {
- public:
-  /**
-   * @brief Constructs an element_hasher object.
-   *
-   * @param nulls Indicates whether to check for nulls
-   * @param seed  The seed to use for the hash function
-   * @param null_hash The hash value to use for nulls
-   */
-  __device__ element_hasher(
-    Nullate nulls,
-    uint32_t seed             = DEFAULT_HASH_SEED,
-    hash_value_type null_hash = cuda::std::numeric_limits<hash_value_type>::max()) noexcept
-    : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
-  {
-  }
-
-  /**
-   * @brief Returns the hash value of the given element.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view const& col,
-                                        size_type row_index) const noexcept
-  {
-    if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
-    return hash_function<T>{_seed}(col.element<T>(row_index));
-  }
-
-  /**
-   * @brief Returns the hash value of the given element.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view const& col,
-                                        size_type row_index) const noexcept
-  {
-    CUDF_UNREACHABLE("Unsupported type in hash.");
-  }
-
-  Nullate _check_nulls;        ///< Whether to check for nulls
-  uint32_t _seed;              ///< The seed to use for hashing
-  hash_value_type _null_hash;  ///< Hash value to use for null elements
-};
-
-/**
- * @brief Computes the hash value of a row in the given table.
- *
- * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-template <template <typename> class hash_function, typename Nullate>
-class device_row_hasher {
-  friend class row_hasher;  ///< Allow row_hasher to access private members.
-
- public:
-  /**
-   * @brief Return the hash value of a row in the given table.
-   *
-   * @param row_index The row index to compute the hash value of
-   * @return The hash value of the row
-   */
-  __device__ auto operator()(size_type row_index) const noexcept
-  {
-    auto it =
-      thrust::make_transform_iterator(_table.begin(), [row_index, this](auto const& column) {
-        return cudf::type_dispatcher<dispatch_storage_type>(
-          column.type(),
-          element_hasher_adapter<hash_function>{_check_nulls, _seed},
-          column,
-          row_index);
-      });
-
-    // Hash each element and combine all the hash values together
-    return detail::accumulate(it, it + _table.num_columns(), _seed, [](auto hash, auto h) {
-      return cudf::hashing::detail::hash_combine(hash, h);
-    });
-  }
-
- private:
-  /**
-   * @brief Computes the hash value of an element in the given column.
-   *
-   * When the column is non-nested, this is a simple wrapper around the element_hasher.
-   * When the column is nested, this uses the element_hasher to hash the shape and values of the
-   * column.
-   */
-  template <template <typename> class hash_fn>
-  class element_hasher_adapter {
-    static constexpr hash_value_type NULL_HASH = cuda::std::numeric_limits<hash_value_type>::max();
-    static constexpr hash_value_type NON_NULL_HASH = 0;
-
-   public:
-    __device__ element_hasher_adapter(Nullate check_nulls, uint32_t seed) noexcept
-      : _element_hasher(check_nulls, seed), _check_nulls(check_nulls)
-    {
-    }
-
-    template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
-    __device__ hash_value_type operator()(column_device_view const& col,
-                                          size_type row_index) const noexcept
-    {
-      return _element_hasher.template operator()<T>(col, row_index);
-    }
-
-    template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
-    __device__ hash_value_type operator()(column_device_view const& col,
-                                          size_type row_index) const noexcept
-    {
-      auto hash                   = hash_value_type{0};
-      column_device_view curr_col = col.slice(row_index, 1);
-      while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) {
-        if (_check_nulls) {
-          auto validity_it = detail::make_validity_iterator<true>(curr_col);
-          hash             = detail::accumulate(
-            validity_it, validity_it + curr_col.size(), hash, [](auto hash, auto is_valid) {
-              return cudf::hashing::detail::hash_combine(hash,
-                                                         is_valid ? NON_NULL_HASH : NULL_HASH);
-            });
-        }
-        if (curr_col.type().id() == type_id::STRUCT) {
-          if (curr_col.num_child_columns() == 0) { return hash; }
-          // Non-empty structs are assumed to be decomposed and contain only one child
-          curr_col = detail::structs_column_device_view(curr_col).get_sliced_child(0);
-        } else if (curr_col.type().id() == type_id::LIST) {
-          auto list_col   = detail::lists_column_device_view(curr_col);
-          auto list_sizes = make_list_size_iterator(list_col);
-          hash            = detail::accumulate(
-            list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
-              return cudf::hashing::detail::hash_combine(hash, hash_fn<size_type>{}(size));
-            });
-          curr_col = list_col.get_sliced_child();
-        }
-      }
-      for (int i = 0; i < curr_col.size(); ++i) {
-        hash = cudf::hashing::detail::hash_combine(
-          hash,
-          type_dispatcher<dispatch_void_if_nested>(curr_col.type(), _element_hasher, curr_col, i));
-      }
-      return hash;
-    }
-
-    element_hasher<hash_fn, Nullate> const _element_hasher;
-    Nullate const _check_nulls;
-  };
-
-  CUDF_HOST_DEVICE device_row_hasher(Nullate check_nulls,
-                                     table_device_view t,
-                                     uint32_t seed = DEFAULT_HASH_SEED) noexcept
-    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
-  {
-  }
-
-  Nullate const _check_nulls;
-  table_device_view const _table;
-  uint32_t const _seed;
-};
-
-// Inject row::equality::preprocessed_table into the row::hash namespace
-// As a result, row::equality::preprocessed_table and row::hash::preprocessed table are the same
-// type and are interchangeable.
-/// Preprocessed table type alias for row hashing operations
-using preprocessed_table = row::equality::preprocessed_table;
-
-/**
- * @brief Computes the hash value of a row in the given table.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/row_operators.cuh instead.
- */
-class row_hasher {
- public:
-  /**
-   * @brief Construct an owning object for hashing the rows of a table
-   *
-   * @param t The table containing rows to hash
-   * @param stream The stream to construct this object on. Not the stream that will be used for
-   * comparisons using this object.
-   */
-  row_hasher(table_view const& t, rmm::cuda_stream_view stream)
-    : d_t(preprocessed_table::create(t, stream))
-  {
-  }
-
-  /**
-   * @brief Construct an owning object for hashing the rows of a table from an existing
-   * preprocessed_table
-   *
-   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
-   * multiple `row_hasher` and `equality::self_comparator` objects.
-   *
-   * @param t A table preprocessed for hashing or equality.
-   */
-  row_hasher(std::shared_ptr<preprocessed_table> t) : d_t{std::move(t)} {}
-
-  /**
-   * @brief Get the hash operator to use on the device
-   *
-   * Returns a unary callable, `F`, with signature `hash_function::hash_value_type F(size_type)`.
-   *
-   * `F(i)` returns the hash of row i.
-   *
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls
-   * @param nullate Indicates if any input column contains nulls
-   * @param seed The seed to use for the hash function
-   * @return A hash operator to use on the device
-   */
-  template <
-    template <typename> class hash_function = cudf::hashing::detail::default_hash,
-    template <template <typename> class, typename> class DeviceRowHasher = device_row_hasher,
-    typename Nullate>
-  DeviceRowHasher<hash_function, Nullate> device_hasher(Nullate nullate = {},
-                                                        uint32_t seed   = DEFAULT_HASH_SEED) const
-  {
-    return DeviceRowHasher<hash_function, Nullate>(nullate, *d_t, seed);
-  }
-
- private:
-  std::shared_ptr<preprocessed_table> d_t;
-};
-
-}  // namespace hash
-
-}  // namespace row
-
-}  // namespace experimental
-}  // namespace CUDF_EXPORT cudf
diff --git a/cpp/include/cudf/table/primitive_row_operators.cuh b/cpp/include/cudf/table/primitive_row_operators.cuh
deleted file mode 100644
index af635e7a07b..00000000000
--- a/cpp/include/cudf/table/primitive_row_operators.cuh
+++ /dev/null
@@ -1,310 +0,0 @@
-/*
- * Copyright (c) 2025, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-/**
- * @file
- * @deprecated This header is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-
-#include <cudf/column/column_device_view.cuh>
-#include <cudf/detail/utilities/assert.cuh>
-#include <cudf/hashing/detail/hash_functions.cuh>
-#include <cudf/hashing/detail/hashing.hpp>
-#include <cudf/table/experimental/row_operators.cuh>
-#include <cudf/table/table_device_view.cuh>
-#include <cudf/table/table_view.hpp>
-#include <cudf/utilities/traits.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
-
-#include <cuda/std/limits>
-#include <cuda/std/type_traits>
-#include <thrust/equal.h>
-
-#include <memory>
-
-namespace CUDF_EXPORT cudf {
-
-/**
- * @brief Checks if a table is compatible with primitive row operations
- *
- * A table is compatible with primitive row operations if it contains exactly one column
- * and that column contains only numeric data types.
- *
- * @param table The table to check for compatibility
- * @return Boolean indicating if the table is compatible with primitive row operations
- *
- * @deprecated This function is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-bool is_primitive_row_op_compatible(cudf::table_view const& table);
-
-namespace row::primitive {
-
-/**
- * @brief Returns `void` if it's not a primitive type
- *
- * @deprecated This type alias is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-template <typename T>
-using primitive_type_t = cuda::std::conditional_t<cudf::is_numeric<T>(), T, void>;
-
-/**
- * @brief Custom dispatcher for primitive types
- *
- * @deprecated This struct is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-template <cudf::type_id Id>
-struct dispatch_primitive_type {
-  using type = primitive_type_t<id_to_type<Id>>;  ///< The underlying type
-};
-
-/**
- * @brief Performs an equality comparison between two elements in two columns.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-class element_equality_comparator {
- public:
-  /**
-   * @brief Compares the specified elements for equality.
-   *
-   * @param lhs The first column
-   * @param rhs The second column
-   * @param lhs_element_index The index of the first element
-   * @param rhs_element_index The index of the second element
-   * @return True if lhs and rhs element are equal
-   */
-  template <typename Element, CUDF_ENABLE_IF(cudf::is_equality_comparable<Element, Element>())>
-  __device__ bool operator()(column_device_view const& lhs,
-                             column_device_view const& rhs,
-                             size_type lhs_element_index,
-                             size_type rhs_element_index) const
-  {
-    return cudf::equality_compare(lhs.element<Element>(lhs_element_index),
-                                  rhs.element<Element>(rhs_element_index));
-  }
-
-  // @cond
-  template <typename Element, CUDF_ENABLE_IF(not cudf::is_equality_comparable<Element, Element>())>
-  __device__ bool operator()(column_device_view const&,
-                             column_device_view const&,
-                             size_type,
-                             size_type) const
-  {
-    CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
-  }
-  // @endcond
-};
-
-/**
- * @brief Performs a relational comparison between two elements in two tables.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-class row_equality_comparator {
- public:
-  /**
-   * @brief Construct a new row equality comparator object
-   *
-   * @param has_nulls Indicates if either input column contains nulls
-   * @param lhs Preprocessed table containing the first element
-   * @param rhs Preprocessed table containing the second element (may be the same as lhs)
-   * @param nulls_are_equal Indicates if two null elements are treated as equivalent
-   */
-  row_equality_comparator(
-    cudf::nullate::DYNAMIC const& has_nulls,
-    std::shared_ptr<cudf::experimental::row::equality::preprocessed_table> lhs,
-    std::shared_ptr<cudf::experimental::row::equality::preprocessed_table> rhs,
-    null_equality nulls_are_equal)
-    : _has_nulls{has_nulls}, _lhs{*lhs}, _rhs{*rhs}, _nulls_are_equal{nulls_are_equal}
-  {
-    CUDF_EXPECTS(_lhs.num_columns() == _rhs.num_columns(), "Mismatched number of columns.");
-  }
-
-  /**
-   * @brief Compares the specified rows for equality.
-   *
-   * @param lhs_row_index The index of the first row to compare (in the lhs table)
-   * @param rhs_row_index The index of the second row to compare (in the rhs table)
-   * @return true if both rows are equal, otherwise false
-   */
-  __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const
-  {
-    auto equal_elements = [this, lhs_row_index, rhs_row_index](column_device_view const& l,
-                                                               column_device_view const& r) {
-      // Handle null comparison for each element
-      if (_has_nulls) {
-        bool const lhs_is_null{l.is_null(lhs_row_index)};
-        bool const rhs_is_null{r.is_null(rhs_row_index)};
-        if (lhs_is_null and rhs_is_null) {
-          return _nulls_are_equal == null_equality::EQUAL;
-        } else if (lhs_is_null != rhs_is_null) {
-          return false;
-        }
-      }
-
-      // Both elements are non-null, compare their values
-      element_equality_comparator comparator;
-      return cudf::type_dispatcher<dispatch_primitive_type>(
-        l.type(), comparator, l, r, lhs_row_index, rhs_row_index);
-    };
-
-    return thrust::equal(thrust::seq, _lhs.begin(), _lhs.end(), _rhs.begin(), equal_elements);
-  }
-
-  /**
-   * @brief Compares the specified rows for equality.
-   *
-   * @param lhs_index The index of the first row to compare (in the lhs table)
-   * @param rhs_index The index of the second row to compare (in the rhs table)
-   * @return Boolean indicating if both rows are equal
-   */
-  __device__ bool operator()(cudf::experimental::row::lhs_index_type lhs_index,
-                             cudf::experimental::row::rhs_index_type rhs_index) const
-  {
-    return (*this)(static_cast<size_type>(lhs_index), static_cast<size_type>(rhs_index));
-  }
-
- private:
-  cudf::nullate::DYNAMIC _has_nulls;
-  table_device_view _lhs;
-  table_device_view _rhs;
-  null_equality _nulls_are_equal;
-};
-
-/**
- * @brief Function object for computing the hash value of a row in a column.
- *
- * @tparam Hash Hash functor to use for hashing elements
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-template <template <typename> class Hash>
-class element_hasher {
- public:
-  /**
-   * @brief Returns the hash value of the given element in the given column.
-   *
-   * @tparam T The type of the element to hash
-   * @param seed The seed value to use for hashing
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(hash_value_type seed,
-                                        column_device_view const& col,
-                                        size_type row_index) const
-  {
-    return Hash<T>{seed}(col.element<T>(row_index));
-  }
-
-  // @cond
-  template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(hash_value_type, column_device_view const&, size_type) const
-  {
-    CUDF_UNREACHABLE("Unsupported type in hash.");
-  }
-  // @endcond
-};
-
-/**
- * @brief Computes the hash value of a row in the given table.
- *
- * @tparam Hash Hash functor to use for hashing elements.
- *
- * @deprecated This class is deprecated in 25.10 and will be removed in 25.12.
- * Users should use cudf/detail/row_operator/primitive_row_operators.cuh instead.
- */
-template <template <typename> class Hash = cudf::hashing::detail::default_hash>
-class row_hasher {
- public:
-  row_hasher() = delete;
-
-  /**
-   * @brief Constructs a row_hasher object with a seed value.
-   *
-   * @param has_nulls Indicates if the input column contains nulls
-   * @param t A table_device_view to hash
-   * @param seed A seed value to use for hashing
-   */
-  row_hasher(cudf::nullate::DYNAMIC const& has_nulls,
-             table_device_view t,
-             hash_value_type seed = DEFAULT_HASH_SEED)
-    : _has_nulls{has_nulls}, _table{t}, _seed{seed}
-  {
-  }
-
-  /**
-   * @brief Constructs a row_hasher object with a seed value.
-   *
-   * @param has_nulls Indicates if the input column contains nulls
-   * @param t Preprocessed table to hash
-   * @param seed A seed value to use for hashing
-   */
-  row_hasher(cudf::nullate::DYNAMIC const& has_nulls,
-             std::shared_ptr<cudf::experimental::row::equality::preprocessed_table> t,
-             hash_value_type seed = DEFAULT_HASH_SEED)
-    : _has_nulls{has_nulls}, _table{*t}, _seed{seed}
-  {
-  }
-
-  /**
-   * @brief Computes the hash value of the row at `row_index` in the `table`
-   *
-   * @param row_index The index of the row in the `table` to hash
-   * @return The hash value of the row at `row_index` in the `table`
-   */
-  __device__ auto operator()(size_type row_index) const
-  {
-    element_hasher<Hash> hasher;
-    // avoid hash combine call if there is only one column
-    auto hash = cuda::std::numeric_limits<hash_value_type>::max();
-    if (!_has_nulls || !_table.column(0).is_null(row_index)) {
-      hash = cudf::type_dispatcher<dispatch_primitive_type>(
-        _table.column(0).type(), hasher, _seed, _table.column(0), row_index);
-    }
-
-    for (size_type i = 1; i < _table.num_columns(); ++i) {
-      if (!(_has_nulls && _table.column(i).is_null(row_index))) {
-        hash = cudf::hashing::detail::hash_combine(
-          hash,
-          cudf::type_dispatcher<dispatch_primitive_type>(
-            _table.column(i).type(), hasher, _seed, _table.column(i), row_index));
-      } else {
-        hash = cudf::hashing::detail::hash_combine(
-          hash, cuda::std::numeric_limits<hash_value_type>::max());
-      }
-    }
-    return hash;
-  }
-
- private:
-  cudf::nullate::DYNAMIC _has_nulls;
-  table_device_view _table;
-  hash_value_type _seed;
-};
-
-}  // namespace row::primitive
-}  // namespace CUDF_EXPORT cudf
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
deleted file mode 100644
index dfcff8977f1..00000000000
--- a/cpp/include/cudf/table/row_operators.cuh
+++ /dev/null
@@ -1,628 +0,0 @@
-/*
- * Copyright (c) 2019-2025, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-/**
- * @file row_operators.cuh
- * @brief DEPRECATED: This header is deprecated as of 25.10 and will be removed in 25.12.
- *
- * Users should migrate to using `cudf::detail::row_operator::row_operators.cuh` directly
- * and access the row operators in the `cudf::detail::row` namespace.
- *
- * @deprecated This header will be removed in cuDF 25.12. Use
- * `cudf::detail::row_operator::row_operators.cuh` instead.
- */
-
-#pragma once
-
-#include <cudf/column/column_device_view.cuh>
-#include <cudf/detail/row_operator/common_utils.cuh>
-#include <cudf/detail/row_operator/row_operators.cuh>
-#include <cudf/detail/utilities/assert.cuh>
-#include <cudf/hashing/detail/hash_functions.cuh>
-#include <cudf/hashing/detail/hashing.hpp>
-#include <cudf/table/table_device_view.cuh>
-#include <cudf/utilities/traits.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
-
-#include <cuda/std/limits>
-#include <thrust/equal.h>
-#include <thrust/execution_policy.h>
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/transform_reduce.h>
-
-namespace CUDF_EXPORT cudf {
-
-// Aliases for backward compatibility - these provide access to common utilities
-// that were previously defined inline or elsewhere but are now in detail namespace
-
-/**
- * @brief Alias for backward compatibility with legacy row operators.
- *
- * @deprecated This alias is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::weak_ordering` instead.
- */
-using weak_ordering = cudf::detail::weak_ordering;
-
-/**
- * @brief Alias for backward compatibility with legacy row operators.
- *
- * @tparam Element The type of the elements being compared
- * @param lhs First element to compare
- * @param rhs Second element to compare
- * @return weak_ordering indicating the relationship between the elements
- * @deprecated This function is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::compare_elements` instead.
- */
-template <typename Element>
-__device__ weak_ordering compare_elements(Element lhs, Element rhs)
-{
-  return cudf::detail::compare_elements(lhs, rhs);
-}
-
-/**
- * @brief Alias for backward compatibility with legacy row operators.
- *
- * @tparam Element The type of the elements being compared
- * @param lhs First element to compare
- * @param rhs Second element to compare
- * @return weak_ordering indicating the relationship between the elements
- * @deprecated This function is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::relational_compare` instead.
- */
-template <typename Element>
-__device__ weak_ordering relational_compare(Element lhs, Element rhs)
-{
-  return cudf::detail::relational_compare(lhs, rhs);
-}
-
-/**
- * @brief Alias for backward compatibility with legacy row operators.
- *
- * @param lhs_is_null Boolean representing if lhs is null
- * @param rhs_is_null Boolean representing if rhs is null
- * @param null_precedence Null order specifying precedence of null values
- * @return weak_ordering indicating the relationship between null in lhs and rhs columns
- * @deprecated This function is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::null_compare` instead.
- */
-inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_order null_precedence)
-{
-  return cudf::detail::null_compare(lhs_is_null, rhs_is_null, null_precedence);
-}
-
-/**
- * @brief Alias for backward compatibility with legacy row operators.
- *
- * @tparam Element The type of the elements being compared
- * @param lhs First element to compare
- * @param rhs Second element to compare
- * @return bool `true` if `lhs` == `rhs` else `false`
- * @deprecated This function is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::equality_compare` instead.
- */
-template <typename Element>
-__device__ bool equality_compare(Element lhs, Element rhs)
-{
-  return cudf::detail::equality_compare(lhs, rhs);
-}
-
-/**
- * @brief Performs an equality comparison between two elements in two columns.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::equality::device_row_comparator::element_comparator` instead.
- *
- * @tparam Nullate A cudf::nullate type describing how to check for nulls.
- */
-template <typename Nullate>
-class element_equality_comparator {
- public:
-  /**
-   * @brief Construct type-dispatched function object for comparing equality
-   * between two elements.
-   *
-   * @note `lhs` and `rhs` may be the same.
-   *
-   * @param has_nulls Indicates if either input column contains nulls.
-   * @param lhs The column containing the first element
-   * @param rhs The column containing the second element (may be the same as lhs)
-   * @param nulls_are_equal Indicates if two null elements are treated as equivalent
-   */
-  __host__ __device__
-  element_equality_comparator(Nullate has_nulls,
-                              column_device_view lhs,
-                              column_device_view rhs,
-                              null_equality nulls_are_equal = null_equality::EQUAL)
-    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal}
-  {
-  }
-
-  /**
-   * @brief Compares the specified elements for equality.
-   *
-   * @param lhs_element_index The index of the first element
-   * @param rhs_element_index The index of the second element
-   * @return True if both lhs and rhs element are both nulls and `nulls_are_equal` is true, or equal
-   */
-  template <typename Element>
-  __device__ bool operator()(size_type lhs_element_index,
-                             size_type rhs_element_index) const noexcept
-    requires(cudf::is_equality_comparable<Element, Element>())
-  {
-    if (nulls) {
-      bool const lhs_is_null{lhs.is_null(lhs_element_index)};
-      bool const rhs_is_null{rhs.is_null(rhs_element_index)};
-      if (lhs_is_null and rhs_is_null) {
-        return nulls_are_equal == null_equality::EQUAL;
-      } else if (lhs_is_null != rhs_is_null) {
-        return false;
-      }
-    }
-
-    return cudf::detail::equality_compare(lhs.element<Element>(lhs_element_index),
-                                          rhs.element<Element>(rhs_element_index));
-  }
-
-  // @cond
-  template <typename Element>
-  __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index)
-    requires(not cudf::is_equality_comparable<Element, Element>())
-  {
-    CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
-  }
-  // @endcond
-
- private:
-  column_device_view lhs;
-  column_device_view rhs;
-  Nullate nulls;
-  null_equality nulls_are_equal;
-};
-
-/**
- * @brief Performs a relational comparison between two elements in two tables.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::equality::device_row_comparator` instead.
- *
- * @tparam Nullate A cudf::nullate type describing how to check for nulls
- */
-template <typename Nullate>
-class row_equality_comparator {
- public:
-  /**
-   * @brief Construct a new row equality comparator object
-   *
-   * @param has_nulls Indicates if either input column contains nulls
-   * @param lhs The column containing the first element
-   * @param rhs The column containing the second element (may be the same as lhs)
-   * @param nulls_are_equal Indicates if two null elements are treated as equivalent
-   */
-  row_equality_comparator(Nullate has_nulls,
-                          table_device_view lhs,
-                          table_device_view rhs,
-                          null_equality nulls_are_equal = null_equality::EQUAL)
-    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal}
-  {
-    CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns.");
-  }
-
-  /**
-   * @brief Compares the specified rows for equality.
-   *
-   * @param lhs_row_index The index of the first row to compare (in the lhs table)
-   * @param rhs_row_index The index of the second row to compare (in the rhs table)
-   * @return true if both rows are equal, otherwise false
-   */
-  __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept
-  {
-    auto equal_elements = [lhs_row_index, rhs_row_index, this](column_device_view l,
-                                                               column_device_view r) {
-      return cudf::type_dispatcher(l.type(),
-                                   element_equality_comparator{nulls, l, r, nulls_are_equal},
-                                   lhs_row_index,
-                                   rhs_row_index);
-    };
-
-    return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements);
-  }
-
- private:
-  table_device_view lhs;
-  table_device_view rhs;
-  Nullate nulls;
-  null_equality nulls_are_equal;
-};
-
-/**
- * @brief Performs a relational comparison between two elements in two columns.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::lexicographic::device_row_comparator::element_comparator` instead.
- *
- * @tparam Nullate A cudf::nullate type describing how to check for nulls.
- */
-template <typename Nullate>
-class element_relational_comparator {
- public:
-  /**
-   * @brief Construct type-dispatched function object for performing a
-   * relational comparison between two elements.
-   *
-   * @note `lhs` and `rhs` may be the same.
-   *
-   * @param lhs The column containing the first element
-   * @param rhs The column containing the second element (may be the same as lhs)
-   * @param has_nulls Indicates if either input column contains nulls.
-   * @param null_precedence Indicates how null values are ordered with other values
-   */
-  __host__ __device__ element_relational_comparator(Nullate has_nulls,
-                                                    column_device_view lhs,
-                                                    column_device_view rhs,
-                                                    null_order null_precedence)
-    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, null_precedence{null_precedence}
-  {
-  }
-
-  /**
-   * @brief Construct type-dispatched function object for performing a relational comparison between
-   * two elements in two columns.
-   *
-   * @param has_nulls Indicates if either input column contains nulls
-   * @param lhs The column containing the first element
-   * @param rhs The column containing the second element (may be the same as lhs)
-   */
-  __host__ __device__ element_relational_comparator(Nullate has_nulls,
-                                                    column_device_view lhs,
-                                                    column_device_view rhs)
-    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}
-  {
-  }
-
-  /**
-   * @brief Performs a relational comparison between the specified elements
-   *
-   * @param lhs_element_index The index of the first element
-   * @param rhs_element_index The index of the second element
-   * @return Indicates the relationship between the elements in
-   * the `lhs` and `rhs` columns.
-   */
-  template <typename Element>
-  __device__ cudf::detail::weak_ordering operator()(size_type lhs_element_index,
-                                                    size_type rhs_element_index) const noexcept
-    requires(cudf::is_relationally_comparable<Element, Element>())
-  {
-    if (nulls) {
-      bool const lhs_is_null{lhs.is_null(lhs_element_index)};
-      bool const rhs_is_null{rhs.is_null(rhs_element_index)};
-
-      if (lhs_is_null or rhs_is_null) {  // at least one is null
-        return cudf::detail::null_compare(lhs_is_null, rhs_is_null, null_precedence);
-      }
-    }
-
-    return cudf::detail::relational_compare(lhs.element<Element>(lhs_element_index),
-                                            rhs.element<Element>(rhs_element_index));
-  }
-
-  // @cond
-  template <typename Element>
-  __device__ cudf::detail::weak_ordering operator()(size_type lhs_element_index,
-                                                    size_type rhs_element_index)
-    requires(not cudf::is_relationally_comparable<Element, Element>())
-  {
-    CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
-  }
-  // @endcond
-
- private:
-  column_device_view lhs;
-  column_device_view rhs;
-  Nullate nulls;
-  null_order null_precedence{};
-};
-
-/**
- * @brief Computes whether one row is lexicographically *less* than another row.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::lexicographic::device_row_comparator` instead.
- *
- * Lexicographic ordering is determined by:
- * - Two rows are compared element by element.
- * - The first mismatching element defines which row is lexicographically less
- * or greater than the other.
- *
- * Lexicographic ordering is exactly equivalent to doing an alphabetical sort of
- * two words, for example, `aac` would be *less* than (or precede) `abb`. The
- * second letter in both words is the first non-equal letter, and `a < b`, thus
- * `aac < abb`.
- *
- * @tparam Nullate A cudf::nullate type describing how to check for nulls.
- */
-template <typename Nullate>
-class row_lexicographic_comparator {
- public:
-  /**
-   * @brief Construct a function object for performing a lexicographic
-   * comparison between the rows of two tables.
-   *
-   * Behavior is undefined if called with incomparable column types.
-   *
-   * @throws cudf::logic_error if `lhs.num_columns() != rhs.num_columns()`
-   *
-   * @param has_nulls Indicates if either input table contains columns with nulls.
-   * @param lhs The first table
-   * @param rhs The second table (may be the same table as `lhs`)
-   * @param column_order Optional, device array the same length as a row that
-   * indicates the desired ascending/descending order of each column in a row.
-   * If `nullptr`, it is assumed all columns are sorted in ascending order.
-   * @param null_precedence Optional, device array the same length as a row
-   * and indicates how null values compare to all other for every column. If
-   * it is nullptr, then null precedence would be `null_order::BEFORE` for all
-   * columns.
-   */
-  row_lexicographic_comparator(Nullate has_nulls,
-                               table_device_view lhs,
-                               table_device_view rhs,
-                               order const* column_order         = nullptr,
-                               null_order const* null_precedence = nullptr)
-    : _lhs{lhs},
-      _rhs{rhs},
-      _nulls{has_nulls},
-      _column_order{column_order},
-      _null_precedence{null_precedence}
-  {
-    CUDF_EXPECTS(_lhs.num_columns() == _rhs.num_columns(), "Mismatched number of columns.");
-  }
-
-  /**
-   * @brief Checks whether the row at `lhs_index` in the `lhs` table compares
-   * lexicographically less than the row at `rhs_index` in the `rhs` table.
-   *
-   * @param lhs_index The index of the row in the `lhs` table to examine
-   * @param rhs_index The index of the row in the `rhs` table to examine
-   * @return `true` if row from the `lhs` table compares less than row in the
-   * `rhs` table
-   */
-  __device__ bool operator()(size_type lhs_index, size_type rhs_index) const noexcept
-  {
-    for (size_type i = 0; i < _lhs.num_columns(); ++i) {
-      bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING);
-
-      null_order null_precedence =
-        _null_precedence == nullptr ? null_order::BEFORE : _null_precedence[i];
-
-      auto comparator =
-        element_relational_comparator{_nulls, _lhs.column(i), _rhs.column(i), null_precedence};
-
-      cudf::detail::weak_ordering state =
-        cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index);
-
-      if (state == cudf::detail::weak_ordering::EQUIVALENT) { continue; }
-
-      return state ==
-             (ascending ? cudf::detail::weak_ordering::LESS : cudf::detail::weak_ordering::GREATER);
-    }
-    return false;
-  }
-
- private:
-  table_device_view _lhs;
-  table_device_view _rhs;
-  Nullate _nulls{};
-  null_order const* _null_precedence{};
-  order const* _column_order{};
-};  // class row_lexicographic_comparator
-
-/**
- * @brief Computes the hash value of an element in the given column.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::hash::element_hasher` instead.
- *
- * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam Nullate A cudf::nullate type describing how to check for nulls.
- */
-template <template <typename> class hash_function, typename Nullate>
-class element_hasher {
- public:
-  /**
-   * @brief Returns the hash value of the given element in the given column.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
-  {
-    if (has_nulls && col.is_null(row_index)) {
-      return cuda::std::numeric_limits<hash_value_type>::max();
-    }
-    return hash_function<T>{}(col.element<T>(row_index));
-  }
-
-  /**
-   * @brief Returns the hash value of the given element in the given column.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
-  {
-    CUDF_UNREACHABLE("Unsupported type in hash.");
-  }
-
-  Nullate has_nulls;  ///< A cudf::nullate type describing how to check for nulls.
-};
-
-/**
- * @brief Function object for computing the hash value of a row in a column.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::hash::element_hasher` instead.
- *
- * @tparam hash_function Hash functor to use for hashing elements
- * @tparam Nullate A cudf::nullate type describing how to check for nulls
- */
-template <template <typename> class hash_function, typename Nullate>
-class element_hasher_with_seed {
- public:
-  /**
-   * @brief Constructs a function object for hashing an element in the given column
-   *
-   * @param has_nulls Indicates if either input column contains nulls
-   * @param seed The seed to use for the hash function
-   */
-  __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed)
-    : _seed{seed}, _has_nulls{has_nulls}
-  {
-  }
-
-  /**
-   * @brief Constructs a function object for hashing an element in the given column
-   *
-   * @param has_nulls Indicates if either input column contains nulls
-   * @param seed The seed to use for the hash function
-   * @param null_hash The hash value to use for null elements
-   */
-  __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed, hash_value_type null_hash)
-    : _seed{seed}, _null_hash{null_hash}, _has_nulls{has_nulls}
-  {
-  }
-
-  /**
-   * @brief Returns the hash value of the given element in the given column.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
-  {
-    if (_has_nulls && col.is_null(row_index)) { return _null_hash; }
-    return hash_function<T>{_seed}(col.element<T>(row_index));
-  }
-
-  /**
-   * @brief Returns the hash value of the given element in the given column.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
-  {
-    CUDF_UNREACHABLE("Unsupported type in hash.");
-  }
-
- private:
-  uint32_t _seed{DEFAULT_HASH_SEED};
-  hash_value_type _null_hash{cuda::std::numeric_limits<hash_value_type>::max()};
-  Nullate _has_nulls;
-};
-
-/**
- * @brief Computes the hash value of a row in the given table.
- *
- * @deprecated This class is deprecated as of 25.10 and will be removed in 25.12.
- * Use `cudf::detail::row::hash::device_row_hasher` instead.
- *
- * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam Nullate A cudf::nullate type describing how to check for nulls.
- */
-template <template <typename> class hash_function, typename Nullate>
-class row_hasher {
- public:
-  row_hasher() = delete;
-
-  /**
-   * @brief Constructs a row_hasher object.
-   *
-   * @param has_nulls Indicates if either input table contains nulls
-   * @param t A table_device_view to hash
-   */
-  CUDF_HOST_DEVICE row_hasher(Nullate has_nulls, table_device_view t)
-    : _table{t}, _has_nulls{has_nulls}
-  {
-  }
-  /**
-   * @brief Constructs a row_hasher object with a seed value.
-   *
-   * @param has_nulls Indicates if either input table contains nulls
-   * @param t A table_device_view to hash
-   * @param seed A seed value to use for hashing
-   */
-  CUDF_HOST_DEVICE row_hasher(Nullate has_nulls, table_device_view t, uint32_t seed)
-    : _table{t}, _seed(seed), _has_nulls{has_nulls}
-  {
-  }
-
-  /**
-   * @brief Computes the hash value of the row at `row_index` in the `table`
-   *
-   * @param row_index The index of the row in the `table` to hash
-   * @return The hash value of the row at `row_index` in the `table`
-   */
-  __device__ auto operator()(size_type row_index) const
-  {
-    // Hash the first column w/ the seed
-    auto const initial_hash = cudf::hashing::detail::hash_combine(
-      hash_value_type{0},
-      type_dispatcher<dispatch_storage_type>(
-        _table.column(0).type(),
-        element_hasher_with_seed<hash_function, Nullate>{_has_nulls, _seed},
-        _table.column(0),
-        row_index));
-
-    // Hashes an element in a column
-    auto hasher = [row_index, this](size_type column_index) {
-      return cudf::type_dispatcher<dispatch_storage_type>(
-        _table.column(column_index).type(),
-        element_hasher<hash_function, Nullate>{_has_nulls},
-        _table.column(column_index),
-        row_index);
-    };
-
-    // Hash each element and combine all the hash values together
-    return thrust::transform_reduce(
-      thrust::seq,
-      // note that this starts at 1 and not 0 now since we already hashed the first column
-      thrust::make_counting_iterator(1),
-      thrust::make_counting_iterator(_table.num_columns()),
-      hasher,
-      initial_hash,
-      [](hash_value_type lhs, hash_value_type rhs) {
-        return cudf::hashing::detail::hash_combine(lhs, rhs);
-      });
-  }
-
- private:
-  table_device_view _table;
-  Nullate _has_nulls;
-  uint32_t _seed{DEFAULT_HASH_SEED};
-};
-
-}  // namespace CUDF_EXPORT cudf
diff --git a/cpp/src/stream_compaction/unique_count_column.cu b/cpp/src/stream_compaction/unique_count_column.cu
index 20e3740a7bf..0acdf217297 100644
--- a/cpp/src/stream_compaction/unique_count_column.cu
+++ b/cpp/src/stream_compaction/unique_count_column.cu
@@ -15,7 +15,6 @@
  */
 
 #include <cudf/column/column_device_view.cuh>
-#include <cudf/column/column_factories.hpp>
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/row_operator/row_operators.cuh>
@@ -28,12 +27,12 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/cmath>
+#include <cuda/std/type_traits>
 #include <thrust/count.h>
 #include <thrust/execution_policy.h>
 #include <thrust/iterator/counting_iterator.h>
 
-#include <cmath>
-
 namespace cudf {
 namespace detail {
 namespace {
@@ -42,17 +41,15 @@ namespace {
  * the row `index` of `column_device_view` is `NaN`.
  */
 struct check_nan {
-  // Check if a value is `NaN` for floating point type columns
   template <typename T>
   __device__ inline bool operator()(column_device_view const& input, size_type index)
-    requires(std::is_floating_point_v<T>)
+    requires(cuda::std::is_floating_point_v<T>)
   {
-    return std::isnan(input.data<T>()[index]);
+    return cuda::std::isnan(input.data<T>()[index]);
   }
-  // Non-floating point type columns can never have `NaN`, so it will always return false.
   template <typename T>
   __device__ inline bool operator()(column_device_view const&, size_type)
-    requires(not std::is_floating_point_v<T>)
+    requires(not cuda::std::is_floating_point_v<T>)
   {
     return false;
   }
@@ -74,11 +71,12 @@ cudf::size_type unique_count(column_view const& input,
   auto input_device_view      = cudf::column_device_view::create(input, stream);
   auto device_view            = *input_device_view;
   auto input_table_view       = table_view{{input}};
-  auto table_ptr              = cudf::table_device_view::create(input_table_view, stream);
-  row_equality_comparator comp(nullate::DYNAMIC{cudf::has_nulls(input_table_view)},
-                               *table_ptr,
-                               *table_ptr,
-                               null_equality::EQUAL);
+
+  auto const comparator = cudf::detail::row::equality::self_comparator{input_table_view, stream};
+  auto const comp       = comparator.equal_to<false>(
+    nullate::DYNAMIC{cudf::has_nulls(input_table_view)},
+    null_equality::EQUAL,
+    cudf::detail::row::equality::nan_equal_physical_equality_comparator{});
 
   return thrust::count_if(
     rmm::exec_policy(stream),
diff --git a/cpp/src/table/primitive_row_operators.cu b/cpp/src/table/primitive_row_operators.cu
deleted file mode 100644
index d879b6d50f7..00000000000
--- a/cpp/src/table/primitive_row_operators.cu
+++ /dev/null
@@ -1,29 +0,0 @@
-/*
- * Copyright (c) 2025, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include <cudf/table/primitive_row_operators.cuh>
-#include <cudf/types.hpp>
-#include <cudf/utilities/traits.hpp>
-
-#include <algorithm>
-
-namespace cudf {
-bool is_primitive_row_op_compatible(cudf::table_view const& table)
-{
-  return std::all_of(
-    table.begin(), table.end(), [](auto const& col) { return cudf::is_numeric(col.type()); });
-}
-}  // namespace cudf
diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu
deleted file mode 100644
index 63ff460dd79..00000000000
--- a/cpp/src/table/row_operators.cu
+++ /dev/null
@@ -1,880 +0,0 @@
-/*
- * Copyright (c) 2022-2025, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include "lists/utilities.hpp"
-
-#include <cudf/column/column.hpp>
-#include <cudf/column/column_factories.hpp>
-#include <cudf/detail/concatenate.hpp>
-#include <cudf/detail/copy.hpp>
-#include <cudf/detail/sorting.hpp>
-#include <cudf/detail/structs/utilities.hpp>
-#include <cudf/detail/utilities/linked_column.hpp>
-#include <cudf/detail/utilities/vector_factories.hpp>
-#include <cudf/lists/lists_column_view.hpp>
-#include <cudf/table/experimental/row_operators.cuh>
-#include <cudf/table/table_view.hpp>
-#include <cudf/utilities/memory_resource.hpp>
-#include <cudf/utilities/type_checks.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
-
-#include <thrust/iterator/transform_iterator.h>
-
-#include <functional>
-
-namespace cudf {
-namespace experimental {
-
-namespace {
-
-/**
- * @brief Removes the offsets of struct column's children
- *
- * @param c The column whose children are to be un-sliced
- * @return Children of `c` with offsets removed
- */
-std::vector<column_view> unslice_children(column_view const& c)
-{
-  if (c.type().id() == type_id::STRUCT) {
-    auto child_it = thrust::make_transform_iterator(c.child_begin(), [](auto const& child) {
-      return column_view(
-        child.type(),
-        child.offset() + child.size(),  // This is hacky, we don't know the actual unsliced size but
-                                        // it is at least offset + size
-        child.head(),
-        child.null_mask(),
-        child.null_count(),
-        0,
-        unslice_children(child));
-    });
-    return {child_it, child_it + c.num_children()};
-  }
-  return {c.child_begin(), c.child_end()};
-};
-
-/**
- * @brief Removes the child column offsets of struct columns in a table.
- *
- * Given a table, this replaces any struct columns with similar struct columns that have their
- * offsets removed from their children. Structs that are children of list columns are not affected.
- *
- */
-table_view remove_struct_child_offsets(table_view table)
-{
-  std::vector<column_view> cols;
-  cols.reserve(table.num_columns());
-  std::transform(table.begin(), table.end(), std::back_inserter(cols), [&](column_view const& c) {
-    return column_view(c.type(),
-                       c.size(),
-                       c.head<uint8_t>(),
-                       c.null_mask(),
-                       c.null_count(),
-                       c.offset(),
-                       unslice_children(c));
-  });
-  return table_view(cols);
-}
-
-/**
- * @brief The enum to specify whether the `decompose_structs` function will process lists columns
- * (at any nested level) or will output them unchanged.
- */
-enum class decompose_lists_column : bool { YES, NO };
-
-/**
- * @brief Decompose all struct columns in a table
- *
- * If a structs column is a tree with N leaves, then this function decomposes the tree into
- * N "linear trees" (branch factor == 1) and prunes common parents. Also returns a vector of
- * per-column `depth`s.
- *
- * A `depth` value is the number of nested levels as parent of the column in the original,
- * non-decomposed table, which are pruned during decomposition.
- *
- * Special handling is needed in the cases of structs column having lists as its first child. In
- * such situations, the function decomposes the tree of N leaves into N+1 linear trees in which the
- * second tree was generated by extracting out leaf of the first tree. This is to make sure there is
- * no structs column having child lists column in the output. Note that structs with lists children
- * in subsequent positions do not require any special treatment because the struct parent will be
- * pruned for all subsequent children.
- *
- * For example, if the original table has a column `Struct<Struct<int, float>, decimal>`,
- *
- *      S1
- *     / \
- *    S2  d
- *   / \
- *  i   f
- *
- * then after decomposition, we get three columns:
- * `Struct<Struct<int>>`, `float`, and `decimal`.
- *
- *  0   2   1  <- depths
- *  S1
- *  |
- *  S2      d
- *  |
- *  i   f
- *
- * The depth of the first column is 0 because it contains all its parent levels, while the depth
- * of the second column is 2 because two of its parent struct levels were pruned.
- *
- * Similarly, a struct column of type `Struct<int, Struct<float, decimal>>` is decomposed as follows
- *
- *     S1
- *    / \
- *   i   S2
- *      / \
- *     f   d
- *
- *  0   1   2  <- depths
- *  S1  S2  d
- *  |   |
- *  i   f
- *
- * In the case of structs column with a lists column as its first child such as
- * `Struct<List<int>, float>`, after decomposition we get three columns `Struct<>`,
- * `List<int>`, and `float`.
- *
- * When list columns are present, depending on the input flag `decompose_lists`, the decomposition
- * can be performed similarly to pure structs but list parent columns are NOT pruned. The list
- * parents are still needed to define the range of elements in the leaf that belong to the same row.
- *
- * For example, if the original table has a column `List<Struct<int, float>>`,
- *
- *    L
- *    |
- *    S
- *   / \
- *  i   f
- *
- * after decomposition, we get two columns
- *
- *  L   L
- *  |   |
- *  S   f
- *  |
- *  i
- *
- * Note that the `decompose_lists` flag should be specified as follow:
- *  - `decompose_lists_column::YES` when preprocessing a table for equality comparison.
- *  - `decompose_lists_column::NO` when preprocessing a table for lexicographic comparison,
- *    since we need to keep all lists columns intact to input into the next preprocessing step.
- *
- * @param table The table whose struct columns to decompose.
- * @param decompose_lists Whether to decompose lists columns
- * @param column_order The per-column order if using output with lexicographic comparison
- * @param null_precedence The per-column null precedence
- * @return A tuple containing a table with all struct columns decomposed, new corresponding column
- *         orders and null precedences and depths of the linearized branches
- */
-auto decompose_structs(table_view table,
-                       decompose_lists_column decompose_lists,
-                       host_span<order const> column_order         = {},
-                       host_span<null_order const> null_precedence = {})
-{
-  auto linked_columns = detail::table_to_linked_columns(table);
-
-  std::vector<column_view> verticalized_columns;
-  std::vector<order> new_column_order;
-  std::vector<null_order> new_null_precedence;
-  std::vector<int> verticalized_col_depths;
-  for (size_t col_idx = 0; col_idx < linked_columns.size(); ++col_idx) {
-    detail::linked_column_view const* col = linked_columns[col_idx].get();
-    if (is_nested(col->type())) {
-      // convert and insert
-      std::vector<std::vector<detail::linked_column_view const*>> flattened;
-      std::function<void(
-        detail::linked_column_view const*, std::vector<detail::linked_column_view const*>*, int)>
-        recursive_child = [&](detail::linked_column_view const* c,
-                              std::vector<detail::linked_column_view const*>* branch,
-                              int depth) {
-          branch->push_back(c);
-          if (decompose_lists == decompose_lists_column::YES && c->type().id() == type_id::LIST) {
-            recursive_child(
-              c->children[lists_column_view::child_column_index].get(), branch, depth + 1);
-          } else if (c->type().id() == type_id::STRUCT) {
-            for (size_t child_idx = 0; child_idx < c->children.size(); ++child_idx) {
-              // When child_idx == 0, we also cut off the current branch if its first child is a
-              // lists column.
-              // In such cases, the last column of the current branch will be `Struct<List,...>` and
-              // it will be modified to empty struct type `Struct<>` later on.
-              if (child_idx > 0 || c->children[0]->type().id() == type_id::LIST) {
-                verticalized_col_depths.push_back(depth + 1);
-                branch = &flattened.emplace_back();
-              }
-              recursive_child(c->children[child_idx].get(), branch, depth + 1);
-            }
-          }
-        };
-      auto& branch = flattened.emplace_back();
-      verticalized_col_depths.push_back(0);
-      recursive_child(col, &branch, 0);
-
-      for (auto const& branch : flattened) {
-        column_view temp_col = *branch.back();
-
-        // Change `Struct<List,...>` into empty struct type `Struct<>`.
-        if (temp_col.type().id() == type_id::STRUCT &&
-            (temp_col.num_children() > 0 && temp_col.child(0).type().id() == type_id::LIST)) {
-          temp_col = column_view(temp_col.type(),
-                                 temp_col.size(),
-                                 temp_col.head(),
-                                 temp_col.null_mask(),
-                                 temp_col.null_count(),
-                                 temp_col.offset(),
-                                 {});
-        }
-
-        for (auto it = branch.crbegin() + 1; it < branch.crend(); ++it) {
-          auto const& prev_col = *(*it);
-          auto children =
-            (prev_col.type().id() == type_id::LIST)
-              ? std::vector<column_view>{*prev_col
-                                            .children[lists_column_view::offsets_column_index],
-                                         temp_col}
-              : std::vector<column_view>{temp_col};
-          temp_col = column_view(prev_col.type(),
-                                 prev_col.size(),
-                                 nullptr,
-                                 prev_col.null_mask(),
-                                 prev_col.null_count(),
-                                 prev_col.offset(),
-                                 std::move(children));
-        }
-        // Traverse upward and include any list columns in the ancestors
-        for (detail::linked_column_view* parent = branch.front()->parent; parent;
-             parent                             = parent->parent) {
-          if (parent->type().id() == type_id::LIST) {
-            // Include this parent
-            temp_col = column_view(
-              parent->type(),
-              parent->size(),
-              nullptr,  // list has no data of its own
-              nullptr,  // If we're going through this then nullmask is already in another branch
-              0,
-              parent->offset(),
-              {*parent->children[lists_column_view::offsets_column_index], temp_col});
-          } else if (parent->type().id() == type_id::STRUCT) {
-            // Replace offset with parent's offset
-            temp_col = column_view(temp_col.type(),
-                                   parent->size(),
-                                   temp_col.head(),
-                                   temp_col.null_mask(),
-                                   temp_col.null_count(),
-                                   parent->offset(),
-                                   {temp_col.child_begin(), temp_col.child_end()});
-          }
-        }
-        verticalized_columns.push_back(temp_col);
-      }
-      if (not column_order.empty()) {
-        new_column_order.insert(new_column_order.end(), flattened.size(), column_order[col_idx]);
-      }
-      if (not null_precedence.empty()) {
-        new_null_precedence.insert(
-          new_null_precedence.end(), flattened.size(), null_precedence[col_idx]);
-      }
-    } else {
-      verticalized_columns.push_back(*col);
-      verticalized_col_depths.push_back(0);
-      if (not column_order.empty()) { new_column_order.push_back(column_order[col_idx]); }
-      if (not null_precedence.empty()) { new_null_precedence.push_back(null_precedence[col_idx]); }
-    }
-  }
-  return std::make_tuple(table_view(verticalized_columns),
-                         std::move(new_column_order),
-                         std::move(new_null_precedence),
-                         std::move(verticalized_col_depths));
-}
-
-/*
- * This helper function generates dremel data for any list-type columns in a
- * table. This data is necessary for lexicographic comparisons.
- */
-auto list_lex_preprocess(table_view const& table, rmm::cuda_stream_view stream)
-{
-  std::vector<detail::dremel_data> dremel_data;
-  auto const num_list_columns = std::count_if(
-    table.begin(), table.end(), [](auto const& col) { return col.type().id() == type_id::LIST; });
-  auto dremel_device_views =
-    cudf::detail::make_empty_host_vector<detail::dremel_device_view>(num_list_columns, stream);
-  for (auto const& col : table) {
-    if (col.type().id() == type_id::LIST) {
-      dremel_data.push_back(detail::get_comparator_data(col, {}, false, stream));
-      dremel_device_views.push_back(dremel_data.back());
-    }
-  }
-  auto d_dremel_device_views = detail::make_device_uvector(
-    dremel_device_views, stream, cudf::get_current_device_resource_ref());
-  return std::make_tuple(std::move(dremel_data), std::move(d_dremel_device_views));
-}
-
-using column_checker_fn_t = std::function<void(column_view const&)>;
-
-/**
- * @brief Check a table for compatibility with lexicographic comparison
- *
- * Checks whether a given table contains columns of non-relationally comparable types.
- */
-void check_lex_compatibility(table_view const& input)
-{
-  // Basically check if there's any LIST of STRUCT or STRUCT of LIST hiding anywhere in the table
-  column_checker_fn_t check_column = [&](column_view const& c) {
-    if (c.type().id() == type_id::LIST) {
-      auto const& list_col = lists_column_view(c);
-      CUDF_EXPECTS(list_col.child().type().id() != type_id::STRUCT,
-                   "Cannot lexicographically compare a table with a LIST of STRUCT column");
-      check_column(list_col.child());
-    } else if (c.type().id() == type_id::STRUCT) {
-      for (auto child = c.child_begin(); child < c.child_end(); ++child) {
-        CUDF_EXPECTS(child->type().id() != type_id::LIST,
-                     "Cannot lexicographically compare a table with a STRUCT of LIST column");
-        check_column(*child);
-      }
-    }
-    if (not is_nested(c.type())) {
-      CUDF_EXPECTS(is_relationally_comparable(c.type()),
-                   "Cannot lexicographic compare a table with a column of type " +
-                     cudf::type_to_name(c.type()));
-    }
-  };
-  for (column_view const& c : input) {
-    check_column(c);
-  }
-}
-
-/**
- * @brief Check a table for compatibility with equality comparison
- *
- * Checks whether a given table contains columns of non-equality comparable types.
- */
-void check_eq_compatibility(table_view const& input)
-{
-  column_checker_fn_t check_column = [&](column_view const& c) {
-    if (not is_nested(c.type())) {
-      CUDF_EXPECTS(is_equality_comparable(c.type()),
-                   "Cannot compare equality for a table with a column of type " +
-                     cudf::type_to_name(c.type()));
-    }
-    for (auto child = c.child_begin(); child < c.child_end(); ++child) {
-      check_column(*child);
-    }
-  };
-  for (column_view const& c : input) {
-    check_column(c);
-  }
-}
-
-void check_shape_compatibility(table_view const& lhs, table_view const& rhs)
-{
-  CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(),
-               "Cannot compare tables with different number of columns");
-  for (size_type i = 0; i < lhs.num_columns(); ++i) {
-    CUDF_EXPECTS(column_types_equivalent(lhs.column(i), rhs.column(i)),
-                 "Cannot compare tables with different column types");
-  }
-}
-
-}  // namespace
-
-namespace row {
-
-namespace lexicographic {
-
-namespace {
-
-/**
- * @brief Replace child of the input lists column by a new child column.
- *
- * If the input is not sliced, just replace the input child by the new_child.
- * Otherwise, we have to generate new offsets and replace both the offsets and the child of the
- * input by the new ones. This is because the new child was generated by ranking and always
- * has zero offset, so it cannot replace the input child if it is sliced.
- *
- * The new generated offsets column needs to be returned and kept alive.
- *
- * @param[in] input The input column_view of type LIST
- * @param[in] new_child A new child column to replace the existing child of the input
- * @param[out] out_cols An array to store the new generated offsets (if applicable)
- * @param[in] stream CUDA stream used for device memory operations and kernel launches
- * @param[in] mr Device memory resource used to allocate the returned column
- * @return An output column_view with child replaced
- */
-auto replace_child(column_view const& input,
-                   column_view const& new_child,
-                   std::vector<std::unique_ptr<column>>& out_cols,
-                   rmm::cuda_stream_view stream,
-                   rmm::device_async_resource_ref mr)
-{
-  auto const make_output = [&input](auto const& offsets_cv, auto const& child_cv) {
-    return column_view{data_type{type_id::LIST},
-                       input.size(),
-                       nullptr,
-                       input.null_mask(),
-                       input.null_count(),
-                       0,
-                       {offsets_cv, child_cv}};
-  };
-
-  if (input.offset() == 0) {
-    return make_output(input.child(lists_column_view::offsets_column_index), new_child);
-  }
-
-  out_cols.emplace_back(
-    cudf::lists::detail::get_normalized_offsets(lists_column_view{input}, stream, mr));
-  return make_output(out_cols.back()->view(), new_child);
-}
-
-/**
- * @brief Compute ranks of the input column.
- *
- * `Dense` rank type must be used for compute ranking of the input for later lexicographic
- * comparison.
- *
- * To understand why, consider: `input = [ [{0, "a"}, {3, "c"}], [{0, "a"}, {2, "b"}] ]`.
- * If first rank is used, `transformed_input = [ [0, 3], [1, 2] ]`. Comparing them will lead
- * to the result row(0) < row(1) which is incorrect.
- * With dense rank, `transformed_input = [ [0, 2], [0, 1] ]`, producing the correct output for
- * lexicographic comparison.
- *
- * In addition, since the input column being ranked is always a nested child column instead of
- * a top-level column, the column order for ranking should be fixed to the same value
- * `order::ASCENDING` in all situations.
- * For example, with the same input above, using column order as `order::ASCENDING` we will have
- * `transformed_input = [ [0, 2], [0, 1] ]`. The output of sorting `transformed_input` will be
- * exactly the same as sorting `input` regardless of the sorting order (ASC or DESC).
- *
- * @param input The input column to compute ranks
- * @param column_null_order The flag indicating how nulls compare to non-null values
- * @param stream CUDA stream used for device memory operations and kernel launches
- * @param mr Device memory resource used to allocate the returned column
- * @return The output rank columns
- */
-auto compute_ranks(column_view const& input,
-                   null_order column_null_order,
-                   rmm::cuda_stream_view stream,
-                   rmm::device_async_resource_ref mr)
-{
-  return cudf::detail::rank(input,
-                            rank_method::DENSE,
-                            order::ASCENDING,
-                            null_policy::EXCLUDE,
-                            column_null_order,
-                            false /*percentage*/,
-                            stream,
-                            mr);
-}
-
-/**
- * @brief Transform any nested lists-of-structs column into lists-of-integers column.
- *
- * For a lists-of-structs column at any nested level, its child structs column will be replaced by a
- * `size_type` column computed as its ranks.
- *
- * If the input column is not lists-of-structs, or does not contain lists-of-structs at any nested
- * level, the input will be passed through without any changes.
- *
- * @param input The input column to transform
- * @param column_null_order The flag indicating how nulls compare to non-null values
- * @param stream CUDA stream used for device memory operations and kernel launches
- * @param mr Device memory resource used to allocate the returned column(s)
- * @return A pair consisting of new column_view representing the transformed input, along with
- *         an array containing its rank column(s) (of `size_type` type) and possibly new list
- *         offsets generated during the transformation process
- */
-std::pair<column_view, std::vector<std::unique_ptr<column>>> transform_lists_of_structs(
-  column_view const& input,
-  null_order column_null_order,
-  rmm::cuda_stream_view stream,
-  rmm::device_async_resource_ref mr)
-{
-  std::vector<std::unique_ptr<column>> out_cols;
-
-  if (input.type().id() == type_id::LIST) {
-    auto const child = cudf::lists_column_view{input}.get_sliced_child(stream);
-
-    // Found a lists-of-structs column.
-    if (child.type().id() == type_id::STRUCT) {
-      out_cols.emplace_back(compute_ranks(child, column_null_order, stream, mr));
-      return {replace_child(input, out_cols.back()->view(), out_cols, stream, mr),
-              std::move(out_cols)};
-    }
-    // Found a lists-of-lists column.
-    else if (child.type().id() == type_id::LIST) {
-      // Recursively call transformation on the child column.
-      auto [new_child, out_cols_child] =
-        transform_lists_of_structs(child, column_null_order, stream, mr);
-
-      // Only transform the current column if its child has been transformed.
-      if (out_cols_child.size() > 0) {
-        out_cols.insert(out_cols.end(),
-                        std::make_move_iterator(out_cols_child.begin()),
-                        std::make_move_iterator(out_cols_child.end()));
-        return {replace_child(input, new_child, out_cols, stream, mr), std::move(out_cols)};
-      }
-      // else: child was not transformed so input is also not transformed.
-    }
-    // else: child is not STRUCT or LIST: no transformation.
-  }
-  // else: lhs.type().id() != type_id::LIST.
-  // In such situations, lhs.type().id() can still be type_id::STRUCT. However, any
-  // structs-of-lists should be decomposed into empty struct type `Struct<>` before being
-  // processed by this function so we do nothing here.
-
-  // Passthrough: nothing changed.
-  return {input, std::move(out_cols)};
-}
-
-/**
- * @brief Transform any nested lists-of-structs column into lists-of-integers column.
- *
- * For a lists-of-structs column at any nested level, its child structs column will be replaced by a
- * `size_type` column computed as its ranks. In addition, equivalent child columns of both input
- * columns (i.e., child columns at the same order, same nested level) will be combined and
- * ranked together.
- *
- * If the input columns are not lists-of-structs, or do not contain lists-of-structs at any nested
- * level, there will not be any changes.
- *
- * @param lhs The input lhs column to transform
- * @param rhs The input rhs column to transform
- * @param column_null_order The flag indicating how nulls compare to non-null values
- * @param stream CUDA stream used for device memory operations and kernel launches
- * @param mr Device memory resource used to allocate the returned column(s)
- * @return A tuple consisting of new column_view(s) representing the transformed input, along with
- *         their rank column(s) (of `size_type` type) and possibly new list offsets generated
- *         during the transformation process
- */
-std::tuple<column_view,
-           column_view,
-           std::vector<std::unique_ptr<column>>,
-           std::vector<std::unique_ptr<column>>>
-transform_lists_of_structs(column_view const& lhs,
-                           column_view const& rhs,
-                           null_order column_null_order,
-                           rmm::cuda_stream_view stream,
-                           rmm::device_async_resource_ref mr)
-{
-  std::vector<std::unique_ptr<column>> out_cols_lhs;
-  std::vector<std::unique_ptr<column>> out_cols_rhs;
-
-  auto const make_output = [&](auto const& new_child_lhs, auto const& new_child_rhs) {
-    return std::tuple{replace_child(lhs, new_child_lhs, out_cols_lhs, stream, mr),
-                      replace_child(rhs, new_child_rhs, out_cols_rhs, stream, mr),
-                      std::move(out_cols_lhs),
-                      std::move(out_cols_rhs)};
-  };
-
-  if (lhs.type().id() == type_id::LIST) {
-    auto const child_lhs = cudf::lists_column_view{lhs}.get_sliced_child(stream);
-    auto const child_rhs = cudf::lists_column_view{rhs}.get_sliced_child(stream);
-
-    // Found a lists-of-structs column.
-    if (child_lhs.type().id() == type_id::STRUCT) {
-      auto const concatenated_children =
-        cudf::detail::concatenate(std::vector<column_view>{child_lhs, child_rhs},
-                                  stream,
-                                  cudf::get_current_device_resource_ref());
-
-      auto const ranks        = compute_ranks(concatenated_children->view(),
-                                       column_null_order,
-                                       stream,
-                                       cudf::get_current_device_resource_ref());
-      auto const ranks_slices = cudf::detail::slice(
-        ranks->view(),
-        {0, child_lhs.size(), child_lhs.size(), child_lhs.size() + child_rhs.size()},
-        stream);
-
-      out_cols_lhs.emplace_back(std::make_unique<column>(ranks_slices.front(), stream, mr));
-      out_cols_rhs.emplace_back(std::make_unique<column>(ranks_slices.back(), stream, mr));
-
-      return make_output(out_cols_lhs.back()->view(), out_cols_rhs.back()->view());
-
-    }
-    // Found a lists-of-lists column.
-    else if (child_lhs.type().id() == type_id::LIST) {
-      // Recursively call transformation on the child column.
-      auto [new_child_lhs, new_child_rhs, out_cols_child_lhs, out_cols_child_rhs] =
-        transform_lists_of_structs(child_lhs, child_rhs, column_null_order, stream, mr);
-
-      // Only transform the current pair of columns if their children have been transformed.
-      if (out_cols_child_lhs.size() > 0 || out_cols_child_rhs.size() > 0) {
-        out_cols_lhs.insert(out_cols_lhs.end(),
-                            std::make_move_iterator(out_cols_child_lhs.begin()),
-                            std::make_move_iterator(out_cols_child_lhs.end()));
-        out_cols_rhs.insert(out_cols_rhs.end(),
-                            std::make_move_iterator(out_cols_child_rhs.begin()),
-                            std::make_move_iterator(out_cols_child_rhs.end()));
-
-        return make_output(new_child_lhs, new_child_rhs);
-      }
-    }
-    // else: child is not STRUCT or LIST: just go to the end of this function, no transformation.
-  }
-  // else: lhs.type().id() != type_id::LIST.
-  // In such situations, lhs.type().id() can still be type_id::STRUCT. However, any
-  // structs-of-lists should be decomposed into empty struct type `Struct<>` before being
-  // processed by this function so we do nothing here.
-
-  // Passthrough: nothing changed.
-  return {lhs, rhs, std::move(out_cols_lhs), std::move(out_cols_rhs)};
-}
-
-}  // namespace
-
-std::shared_ptr<preprocessed_table> preprocessed_table::create(
-  table_view const& preprocessed_input,
-  std::vector<int>&& verticalized_col_depths,
-  std::vector<std::unique_ptr<column>>&& transformed_columns,
-  host_span<order const> column_order,
-  host_span<null_order const> null_precedence,
-  bool has_ranked_children,
-  rmm::cuda_stream_view stream)
-{
-  check_lex_compatibility(preprocessed_input);
-
-  auto d_table        = table_device_view::create(preprocessed_input, stream);
-  auto d_column_order = detail::make_device_uvector_async(
-    column_order, stream, cudf::get_current_device_resource_ref());
-  auto d_null_precedence = detail::make_device_uvector_async(
-    null_precedence, stream, cudf::get_current_device_resource_ref());
-  auto d_depths = detail::make_device_uvector_async(
-    verticalized_col_depths, stream, cudf::get_current_device_resource_ref());
-
-  if (detail::has_nested_columns(preprocessed_input)) {
-    auto [dremel_data, d_dremel_device_view] = list_lex_preprocess(preprocessed_input, stream);
-    return std::shared_ptr<preprocessed_table>(
-      new preprocessed_table(std::move(d_table),
-                             std::move(d_column_order),
-                             std::move(d_null_precedence),
-                             std::move(d_depths),
-                             std::move(dremel_data),
-                             std::move(d_dremel_device_view),
-                             std::move(transformed_columns),
-                             has_ranked_children));
-  } else {
-    return std::shared_ptr<preprocessed_table>(
-      new preprocessed_table(std::move(d_table),
-                             std::move(d_column_order),
-                             std::move(d_null_precedence),
-                             std::move(d_depths),
-                             std::move(transformed_columns),
-                             has_ranked_children));
-  }
-}
-
-std::shared_ptr<preprocessed_table> preprocessed_table::create(
-  table_view const& input,
-  host_span<order const> column_order,
-  host_span<null_order const> null_precedence,
-  rmm::cuda_stream_view stream)
-{
-  auto [decomposed_input, new_column_order, new_null_precedence, verticalized_col_depths] =
-    decompose_structs(input, decompose_lists_column::NO, column_order, null_precedence);
-
-  // Transform any (nested) lists-of-structs column into lists-of-integers column.
-  std::vector<std::unique_ptr<column>> transformed_columns;
-  auto const transformed_input =
-    [&, &decomposed_input = decomposed_input, &new_null_precedence = new_null_precedence] {
-      std::vector<column_view> transformed_cvs;
-
-      for (size_type col_idx = 0; col_idx < decomposed_input.num_columns(); ++col_idx) {
-        auto const& lhs_col = decomposed_input.column(col_idx);
-
-        auto [transformed, curr_out_cols] = transform_lists_of_structs(
-          lhs_col,
-          null_precedence.empty() ? null_order::BEFORE : new_null_precedence[col_idx],
-          stream,
-          cudf::get_current_device_resource_ref());
-
-        transformed_cvs.emplace_back(std::move(transformed));
-        transformed_columns.insert(transformed_columns.end(),
-                                   std::make_move_iterator(curr_out_cols.begin()),
-                                   std::make_move_iterator(curr_out_cols.end()));
-      }
-
-      return table_view{transformed_cvs};
-    }();
-
-  auto const has_ranked_children = !transformed_columns.empty();
-  return create(transformed_input,
-                std::move(verticalized_col_depths),
-                std::move(transformed_columns),
-                new_column_order,
-                new_null_precedence,
-                has_ranked_children,
-                stream);
-}
-
-std::pair<std::shared_ptr<preprocessed_table>, std::shared_ptr<preprocessed_table>>
-preprocessed_table::create(table_view const& lhs,
-                           table_view const& rhs,
-                           host_span<order const> column_order,
-                           host_span<null_order const> null_precedence,
-                           rmm::cuda_stream_view stream)
-{
-  check_shape_compatibility(lhs, rhs);
-
-  auto [decomposed_lhs,
-        new_column_order_lhs,
-        new_null_precedence_lhs,
-        verticalized_col_depths_lhs] =
-    decompose_structs(lhs, decompose_lists_column::NO, column_order, null_precedence);
-
-  // Unused variables are new column order and null order for rhs, which are the same as for lhs
-  // so we don't need them.
-  [[maybe_unused]] auto [decomposed_rhs, unused0, unused1, verticalized_col_depths_rhs] =
-    decompose_structs(rhs, decompose_lists_column::NO, column_order, null_precedence);
-
-  // Transform any (nested) lists-of-structs column into lists-of-integers column.
-  std::vector<std::unique_ptr<column>> transformed_columns_lhs;
-  std::vector<std::unique_ptr<column>> transformed_columns_rhs;
-  auto const [transformed_lhs,
-              transformed_rhs] = [&,
-                                  &decomposed_lhs          = decomposed_lhs,
-                                  &decomposed_rhs          = decomposed_rhs,
-                                  &new_null_precedence_lhs = new_null_precedence_lhs] {
-    std::vector<column_view> transformed_lhs_cvs;
-    std::vector<column_view> transformed_rhs_cvs;
-
-    for (size_type col_idx = 0; col_idx < decomposed_lhs.num_columns(); ++col_idx) {
-      auto const& lhs_col = decomposed_lhs.column(col_idx);
-      auto const& rhs_col = decomposed_rhs.column(col_idx);
-
-      auto [transformed_lhs, transformed_rhs, curr_out_cols_lhs, curr_out_cols_rhs] =
-        transform_lists_of_structs(
-          lhs_col,
-          rhs_col,
-          null_precedence.empty() ? null_order::BEFORE : null_precedence[col_idx],
-          stream,
-          cudf::get_current_device_resource_ref());
-
-      transformed_lhs_cvs.emplace_back(std::move(transformed_lhs));
-      transformed_rhs_cvs.emplace_back(std::move(transformed_rhs));
-      transformed_columns_lhs.insert(transformed_columns_lhs.end(),
-                                     std::make_move_iterator(curr_out_cols_lhs.begin()),
-                                     std::make_move_iterator(curr_out_cols_lhs.end()));
-      transformed_columns_rhs.insert(transformed_columns_rhs.end(),
-                                     std::make_move_iterator(curr_out_cols_rhs.begin()),
-                                     std::make_move_iterator(curr_out_cols_rhs.end()));
-    }
-
-    return std::pair{table_view{transformed_lhs_cvs}, table_view{transformed_rhs_cvs}};
-  }();
-
-  // This should be the same for both lhs and rhs but not all the time, such as when one table
-  // has 0 rows while the other has >0 rows. So we check separately for each of them.
-  auto const has_ranked_children_lhs = !transformed_columns_lhs.empty();
-  auto const has_ranked_children_rhs = !transformed_columns_rhs.empty();
-
-  return {create(transformed_lhs,
-                 std::move(verticalized_col_depths_lhs),
-                 std::move(transformed_columns_lhs),
-                 new_column_order_lhs,
-                 new_null_precedence_lhs,
-                 has_ranked_children_lhs,
-                 stream),
-          create(transformed_rhs,
-                 std::move(verticalized_col_depths_rhs),
-                 std::move(transformed_columns_rhs),
-                 new_column_order_lhs,
-                 new_null_precedence_lhs,
-                 has_ranked_children_rhs,
-                 stream)};
-}
-
-preprocessed_table::preprocessed_table(
-  table_device_view_owner&& table,
-  rmm::device_uvector<order>&& column_order,
-  rmm::device_uvector<null_order>&& null_precedence,
-  rmm::device_uvector<size_type>&& depths,
-  std::vector<detail::dremel_data>&& dremel_data,
-  rmm::device_uvector<detail::dremel_device_view>&& dremel_device_views,
-  std::vector<std::unique_ptr<column>>&& transformed_columns,
-  bool has_ranked_children)
-  : _t(std::move(table)),
-    _column_order(std::move(column_order)),
-    _null_precedence(std::move(null_precedence)),
-    _depths(std::move(depths)),
-    _dremel_data(std::move(dremel_data)),
-    _dremel_device_views(std::move(dremel_device_views)),
-    _transformed_columns(std::move(transformed_columns)),
-    _has_ranked_children(has_ranked_children)
-{
-}
-
-preprocessed_table::preprocessed_table(table_device_view_owner&& table,
-                                       rmm::device_uvector<order>&& column_order,
-                                       rmm::device_uvector<null_order>&& null_precedence,
-                                       rmm::device_uvector<size_type>&& depths,
-                                       std::vector<std::unique_ptr<column>>&& transformed_columns,
-                                       bool has_ranked_children)
-  : _t(std::move(table)),
-    _column_order(std::move(column_order)),
-    _null_precedence(std::move(null_precedence)),
-    _depths(std::move(depths)),
-    _dremel_data{},
-    _dremel_device_views{},
-    _transformed_columns(std::move(transformed_columns)),
-    _has_ranked_children(has_ranked_children)
-{
-}
-
-two_table_comparator::two_table_comparator(table_view const& left,
-                                           table_view const& right,
-                                           host_span<order const> column_order,
-                                           host_span<null_order const> null_precedence,
-                                           rmm::cuda_stream_view stream)
-{
-  std::tie(d_left_table, d_right_table) =
-    preprocessed_table::create(left, right, column_order, null_precedence, stream);
-}
-
-}  // namespace lexicographic
-
-namespace equality {
-
-std::shared_ptr<preprocessed_table> preprocessed_table::create(table_view const& t,
-                                                               rmm::cuda_stream_view stream)
-{
-  check_eq_compatibility(t);
-
-  auto [null_pushed_table, nullable_data] =
-    structs::detail::push_down_nulls(t, stream, cudf::get_current_device_resource_ref());
-  auto struct_offset_removed_table = remove_struct_child_offsets(null_pushed_table);
-  auto verticalized_t =
-    std::get<0>(decompose_structs(struct_offset_removed_table, decompose_lists_column::YES));
-
-  auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream));
-  return std::shared_ptr<preprocessed_table>(new preprocessed_table(
-    std::move(d_t), std::move(nullable_data.new_null_masks), std::move(nullable_data.new_columns)));
-}
-
-two_table_comparator::two_table_comparator(table_view const& left,
-                                           table_view const& right,
-                                           rmm::cuda_stream_view stream)
-  : d_left_table{preprocessed_table::create(left, stream)},
-    d_right_table{preprocessed_table::create(right, stream)}
-{
-  check_shape_compatibility(left, right);
-}
-
-}  // namespace equality
-
-}  // namespace row
-}  // namespace experimental
-}  // namespace cudf
diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index ba6d62016af..def973e7b9f 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -486,14 +486,17 @@ ConfigureTest(
 
 # ##################################################################################################
 # * table tests -----------------------------------------------------------------------------------
+ConfigureTest(TABLE_TEST table/table_tests.cpp table/table_view_tests.cu)
+
+# ##################################################################################################
+# * row operator tests ----------------------------------------------------------------------------
 ConfigureTest(
-  TABLE_TEST
-  table/table_tests.cpp
-  table/table_view_tests.cu
-  table/row_operators_tests.cpp
-  table/experimental_row_operator_tests.cu
-  table/row_operator_tests_utilities.cu
-  table/row_operator_tests_utilities2.cu
+  ROW_OPERATOR_TEST
+  row_operator/row_operator_tests.cu
+  row_operator/nan_tests.cpp
+  row_operator/self_comparison_utilities.cu
+  row_operator/two_table_comparison_utilities.cu
+  row_operator/two_table_equality_utilities.cu
 )
 
 # ##################################################################################################
diff --git a/cpp/tests/table/row_operators_tests.cpp b/cpp/tests/row_operator/nan_tests.cpp
similarity index 100%
rename from cpp/tests/table/row_operators_tests.cpp
rename to cpp/tests/row_operator/nan_tests.cpp
diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu
similarity index 100%
rename from cpp/tests/table/experimental_row_operator_tests.cu
rename to cpp/tests/row_operator/row_operator_tests.cu
diff --git a/cpp/tests/table/row_operator_tests_utilities.hpp b/cpp/tests/row_operator/row_operator_tests_utilities.hpp
similarity index 100%
rename from cpp/tests/table/row_operator_tests_utilities.hpp
rename to cpp/tests/row_operator/row_operator_tests_utilities.hpp
diff --git a/cpp/tests/table/row_operator_tests_utilities2.cu b/cpp/tests/row_operator/self_comparison_utilities.cu
similarity index 97%
rename from cpp/tests/table/row_operator_tests_utilities2.cu
rename to cpp/tests/row_operator/self_comparison_utilities.cu
index d85ba0385be..98be10f64e9 100644
--- a/cpp/tests/table/row_operator_tests_utilities2.cu
+++ b/cpp/tests/row_operator/self_comparison_utilities.cu
@@ -25,7 +25,7 @@
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/transform.h>
 
-// Including this declaration/defintion in row_operator_tests_utilities.cu causes
+// Including this declaration/defintion in two_table_comparison_utilities.cu causes
 // the nvcc compiler to segfault when built with the debug (-g) flag.
 
 template <typename PhysicalElementComparator>
diff --git a/cpp/tests/table/row_operator_tests_utilities.cu b/cpp/tests/row_operator/two_table_comparison_utilities.cu
similarity index 69%
rename from cpp/tests/table/row_operator_tests_utilities.cu
rename to cpp/tests/row_operator/two_table_comparison_utilities.cu
index 4549cd361c0..fa173768191 100644
--- a/cpp/tests/table/row_operator_tests_utilities.cu
+++ b/cpp/tests/row_operator/two_table_comparison_utilities.cu
@@ -111,54 +111,3 @@ template std::unique_ptr<cudf::column> sorted_order<sorting_comparator_t>(
   bool has_nested,
   sorting_comparator_t comparator,
   rmm::cuda_stream_view stream);
-
-template <typename PhysicalElementComparator>
-std::unique_ptr<cudf::column> two_table_equality(cudf::table_view lhs,
-                                                 cudf::table_view rhs,
-                                                 std::vector<cudf::order> const& column_order,
-                                                 PhysicalElementComparator comparator)
-{
-  rmm::cuda_stream_view stream{cudf::get_default_stream()};
-
-  auto const table_comparator = cudf::detail::row::equality::two_table_comparator{lhs, rhs, stream};
-
-  auto const lhs_it = cudf::detail::row::lhs_iterator(0);
-  auto const rhs_it = cudf::detail::row::rhs_iterator(0);
-
-  auto output = cudf::make_numeric_column(
-    cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED);
-
-  if (cudf::has_nested_columns(lhs) or cudf::has_nested_columns(rhs)) {
-    auto const equal_comparator =
-      table_comparator.equal_to<true>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
-
-    thrust::transform(rmm::exec_policy(stream),
-                      lhs_it,
-                      lhs_it + lhs.num_rows(),
-                      rhs_it,
-                      output->mutable_view().data<bool>(),
-                      equal_comparator);
-  } else {
-    auto const equal_comparator =
-      table_comparator.equal_to<false>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
-
-    thrust::transform(rmm::exec_policy(stream),
-                      lhs_it,
-                      lhs_it + lhs.num_rows(),
-                      rhs_it,
-                      output->mutable_view().data<bool>(),
-                      equal_comparator);
-  }
-  return output;
-}
-
-template std::unique_ptr<cudf::column> two_table_equality<physical_equality_t>(
-  cudf::table_view lhs,
-  cudf::table_view rhs,
-  std::vector<cudf::order> const& column_order,
-  physical_equality_t comparator);
-template std::unique_ptr<cudf::column> two_table_equality<nan_equality_t>(
-  cudf::table_view lhs,
-  cudf::table_view rhs,
-  std::vector<cudf::order> const& column_order,
-  nan_equality_t comparator);
diff --git a/cpp/tests/row_operator/two_table_equality_utilities.cu b/cpp/tests/row_operator/two_table_equality_utilities.cu
new file mode 100644
index 00000000000..86281167bba
--- /dev/null
+++ b/cpp/tests/row_operator/two_table_equality_utilities.cu
@@ -0,0 +1,76 @@
+/*
+ * Copyright (c) 2025, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "row_operator_tests_utilities.hpp"
+
+#include <cudf/column/column_factories.hpp>
+#include <cudf/column/column_view.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/transform.h>
+
+template <typename PhysicalElementComparator>
+std::unique_ptr<cudf::column> two_table_equality(cudf::table_view lhs,
+                                                 cudf::table_view rhs,
+                                                 std::vector<cudf::order> const& column_order,
+                                                 PhysicalElementComparator comparator)
+{
+  rmm::cuda_stream_view stream{cudf::get_default_stream()};
+
+  auto const table_comparator = cudf::detail::row::equality::two_table_comparator{lhs, rhs, stream};
+
+  auto const lhs_it = cudf::detail::row::lhs_iterator(0);
+  auto const rhs_it = cudf::detail::row::rhs_iterator(0);
+
+  auto output = cudf::make_numeric_column(
+    cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED);
+
+  if (cudf::has_nested_columns(lhs) or cudf::has_nested_columns(rhs)) {
+    auto const equal_comparator =
+      table_comparator.equal_to<true>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
+
+    thrust::transform(rmm::exec_policy(stream),
+                      lhs_it,
+                      lhs_it + lhs.num_rows(),
+                      rhs_it,
+                      output->mutable_view().data<bool>(),
+                      equal_comparator);
+  } else {
+    auto const equal_comparator =
+      table_comparator.equal_to<false>(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator);
+
+    thrust::transform(rmm::exec_policy(stream),
+                      lhs_it,
+                      lhs_it + lhs.num_rows(),
+                      rhs_it,
+                      output->mutable_view().data<bool>(),
+                      equal_comparator);
+  }
+  return output;
+}
+
+template std::unique_ptr<cudf::column> two_table_equality<physical_equality_t>(
+  cudf::table_view lhs,
+  cudf::table_view rhs,
+  std::vector<cudf::order> const& column_order,
+  physical_equality_t comparator);
+template std::unique_ptr<cudf::column> two_table_equality<nan_equality_t>(
+  cudf::table_view lhs,
+  cudf::table_view rhs,
+  std::vector<cudf::order> const& column_order,
+  nan_equality_t comparator);
diff --git a/cpp/tests/table/table_view_tests.cu b/cpp/tests/table/table_view_tests.cu
index ffa3fa466ef..fbd08e2fd07 100644
--- a/cpp/tests/table/table_view_tests.cu
+++ b/cpp/tests/table/table_view_tests.cu
@@ -20,8 +20,8 @@
 #include <cudf_test/type_lists.hpp>
 
 #include <cudf/column/column_view.hpp>
+#include <cudf/detail/row_operator/row_operators.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>
-#include <cudf/table/row_operators.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/table/table_view.hpp>
 #include <cudf/utilities/default_stream.hpp>
@@ -45,20 +45,17 @@ void row_comparison(cudf::table_view input1,
 {
   rmm::cuda_stream_view stream{cudf::get_default_stream()};
 
-  auto device_table_1 = cudf::table_device_view::create(input1, stream);
-  auto device_table_2 = cudf::table_device_view::create(input2, stream);
-  auto d_column_order = cudf::detail::make_device_uvector(
-    column_order, cudf::get_default_stream(), cudf::get_current_device_resource_ref());
-
-  auto comparator = cudf::row_lexicographic_comparator(
-    cudf::nullate::NO{}, *device_table_1, *device_table_2, d_column_order.data());
+  auto const comparator = cudf::detail::row::lexicographic::two_table_comparator{
+    input1, input2, column_order, {}, stream};
+  auto const lhs_it = cudf::detail::row::lhs_iterator(0);
+  auto const rhs_it = cudf::detail::row::rhs_iterator(0);
 
   thrust::transform(rmm::exec_policy(stream),
-                    thrust::make_counting_iterator(0),
-                    thrust::make_counting_iterator(input1.num_rows()),
-                    thrust::make_counting_iterator(0),
+                    lhs_it,
+                    lhs_it + input1.num_rows(),
+                    rhs_it,
                     output.data<int8_t>(),
-                    comparator);
+                    comparator.less<false>(cudf::nullate::NO{}));
 }
 
 TEST_F(TableViewTest, EmptyColumnedTable)