Revert changes to non-experimental row operators.

cpp/include/cudf/table/row_operators.cuh

@@ -20,7 +20,6 @@
 #include <cudf/detail/hashing.hpp>
 #include <cudf/detail/utilities/assert.cuh>
 #include <cudf/detail/utilities/hash_functions.cuh>
-#include <cudf/detail/utilities/strong_index.hpp>
 #include <cudf/sorting.hpp>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/utilities/traits.hpp>
@@ -199,37 +198,30 @@ class element_equality_comparator {
    */
   template <typename Element,
             std::enable_if_t<cudf::is_equality_comparable<Element, Element>()>* = nullptr>
-  __device__ bool operator()(cudf::lhs_index_type lhs_element_index,
-                             cudf::rhs_index_type rhs_element_index) const noexcept
+  __device__ bool operator()(size_type lhs_element_index,
+                             size_type rhs_element_index) const noexcept
   {
     if (nulls) {
-      bool const lhs_is_null{lhs.is_null(lhs_element_index.value())};
-      bool const rhs_is_null{rhs.is_null(rhs_element_index.value())};
+      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 equality_compare(lhs.element<Element>(lhs_element_index.value()),
-                            rhs.element<Element>(rhs_element_index.value()));
+    return equality_compare(lhs.element<Element>(lhs_element_index),
+                            rhs.element<Element>(rhs_element_index));
   }
 
   template <typename Element,
             std::enable_if_t<not cudf::is_equality_comparable<Element, Element>()>* = nullptr>
-  __device__ bool operator()(cudf::lhs_index_type lhs_element_index,
-                             cudf::rhs_index_type rhs_element_index) const noexcept
+  __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index)
   {
     CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
   }
 
-  __device__ bool operator()(cudf::rhs_index_type rhs_element_index,
-                             cudf::lhs_index_type lhs_element_index) const noexcept
-  {
-    return operator()(lhs_element_index, rhs_element_index);
-  }
-
  private:
   column_device_view lhs;
   column_device_view rhs;
@@ -254,8 +246,8 @@ class row_equality_comparator {
     auto equal_elements = [=](column_device_view l, column_device_view r) {
       return cudf::type_dispatcher(l.type(),
                                    element_equality_comparator{nulls, l, r, nulls_are_equal},
-                                   cudf::lhs_index_type(lhs_row_index),
-                                   cudf::rhs_index_type(rhs_row_index));
+                                   lhs_row_index,
+                                   rhs_row_index);
     };
 
     return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements);

cpp/src/groupby/sort/group_nunique.cu

@@ -17,7 +17,6 @@
 #include <cudf/aggregation.hpp>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/iterator.cuh>
-#include <cudf/detail/utilities/strong_index.hpp>
 #include <cudf/table/row_operators.cuh>
 #include <cudf/types.hpp>
 #include <cudf/utilities/span.hpp>
@@ -63,11 +62,9 @@ struct nunique_functor {
          group_labels  = group_labels.data()] __device__(auto i) -> size_type {
           bool is_input_countable =
             (null_handling == null_policy::INCLUDE || v.is_valid_nocheck(i));
-          bool is_unique =
-            is_input_countable &&
-            (group_offsets[group_labels[i]] == i ||  // first element or
-             (not equal.operator()<T>(cudf::lhs_index_type(i),
-                                      cudf::rhs_index_type(i - 1))));  // new unique value in sorted
+          bool is_unique = is_input_countable &&
+                           (group_offsets[group_labels[i]] == i ||  // first element or
+                            (not equal.operator()<T>(i, i - 1)));   // new unique value in sorted
           return static_cast<size_type>(is_unique);
         });
 
@@ -85,10 +82,8 @@ struct nunique_functor {
          equal,
          group_offsets = group_offsets.data(),
          group_labels  = group_labels.data()] __device__(auto i) -> size_type {
-          bool is_unique =
-            group_offsets[group_labels[i]] == i ||  // first element or
-            (not equal.operator()<T>(cudf::lhs_index_type(i),
-                                     cudf::rhs_index_type(i - 1)));  // new unique value in sorted
+          bool is_unique = group_offsets[group_labels[i]] == i ||  // first element or
+                           (not equal.operator()<T>(i, i - 1));    // new unique value in sorted
           return static_cast<size_type>(is_unique);
         });
       thrust::reduce_by_key(rmm::exec_policy(stream),

cpp/src/transform/one_hot_encode.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,7 +19,6 @@
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
-#include <cudf/detail/utilities/strong_index.hpp>
 #include <cudf/table/row_operators.cuh>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
@@ -48,8 +47,8 @@ struct one_hot_encode_functor {
 
   bool __device__ operator()(size_type i)
   {
-    cudf::lhs_index_type const element_index(i % _input_size);
-    cudf::rhs_index_type const category_index(i / _input_size);
+    size_type const element_index  = i % _input_size;
+    size_type const category_index = i / _input_size;
     return _equality_comparator.template operator()<InputType>(element_index, category_index);
   }