Remove is_relationally_comparable for table device views

cpp/include/cudf/table/row_operators.cuh

@@ -358,30 +358,32 @@ class row_lexicographic_comparator {
    * @throws cudf::logic_error if `lhs.num_columns() != rhs.num_columns()`
    * @throws cudf::logic_error if column types of `lhs` and `rhs` are not comparable.
    *
+   * @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 has_nulls Indicates if either input table contains columns with nulls.
    * @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.
+   * @param stream CUDA stream used for device memory operations and kernel launches
    */
   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)
+                               null_order const* null_precedence = nullptr,
+                               rmm::cuda_stream_view stream      = rmm::cuda_stream_default)
     : _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.");
-    CUDF_EXPECTS(detail::is_relationally_comparable(_lhs, _rhs),
+    CUDF_EXPECTS(detail::is_relationally_comparable(_lhs, _rhs, stream),
                  "Attempted to compare elements of uncomparable types.");
   }
 

cpp/include/cudf/table/table_device_view.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -151,9 +151,26 @@ auto contiguous_copy_column_device_views(HostTableView source_view, rmm::cuda_st
 }
 
 namespace detail {
+/**
+ * @brief Indicates whether respective columns in device tables are relationally comparable.
+ *
+ * @param lhs The first table
+ * @param rhs The second table (may be the same table as `lhs`)
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @return true all of respective columns on `lhs` and 'rhs` tables are comparable.
+ * @return false any of respective columns on `lhs` and 'rhs` tables are not comparable.
+ */
+template <typename TableDeviceView>
+bool is_relationally_comparable(TableDeviceView const& lhs,
+                                TableDeviceView const& rhs,
+                                rmm::cuda_stream_view stream);
+
 extern template bool is_relationally_comparable<table_device_view>(table_device_view const& lhs,
-                                                                   table_device_view const& rhs);
+                                                                   table_device_view const& rhs,
+                                                                   rmm::cuda_stream_view stream);
 extern template bool is_relationally_comparable<mutable_table_device_view>(
-  mutable_table_device_view const& lhs, mutable_table_device_view const& rhs);
+  mutable_table_device_view const& lhs,
+  mutable_table_device_view const& rhs,
+  rmm::cuda_stream_view stream);
 }  // namespace detail
 }  // namespace cudf

cpp/src/groupby/sort/group_scan_util.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -201,7 +201,7 @@ struct group_scan_functor<K,
                                   thrust::make_counting_iterator<size_type>(0),
                                   gather_map.begin(),
                                   thrust::equal_to{},
-                                  binop_generator.binop());
+                                  binop_generator.binop(stream));
 
     //
     // Gather the children elements of the prefix min/max struct elements first.

cpp/src/groupby/sort/group_single_pass_reduction_util.cuh

@@ -273,7 +273,7 @@ struct group_reduction_functor<
     auto const result_begin = result->mutable_view().template begin<ResultType>();
     auto const binop_generator =
       cudf::reduction::detail::comparison_binop_generator::create<K>(values, stream);
-    do_reduction(count_iter, result_begin, binop_generator.binop());
+    do_reduction(count_iter, result_begin, binop_generator.binop(stream));
 
     if (values.has_nulls()) {
       // Generate bitmask for the output by segmented reduction of the input bitmask.

cpp/src/reductions/scan/scan_inclusive.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -164,7 +164,7 @@ struct scan_functor<Op, cudf::struct_view> {
                            thrust::counting_iterator<size_type>(0),
                            thrust::counting_iterator<size_type>(input.size()),
                            gather_map.begin(),
-                           binop_generator.binop());
+                           binop_generator.binop(stream));
 
     // Gather the children columns of the input column. Must use `get_sliced_child` to properly
     // handle input in case it is a sliced view.

cpp/src/reductions/simple.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -300,7 +300,7 @@ struct same_element_type_dispatcher {
                                            thrust::make_counting_iterator(0),
                                            thrust::make_counting_iterator(input.size()),
                                            size_type{0},
-                                           binop_generator.binop());
+                                           binop_generator.binop(stream));
 
     return cudf::detail::get_element(input, minmax_idx, stream, mr);
   }

cpp/src/reductions/struct_minmax_util.cuh

@@ -40,9 +40,10 @@ struct row_arg_minmax_fn {
   row_arg_minmax_fn(table_device_view const& table,
                     bool has_nulls,
                     null_order const* null_precedence,
-                    bool const arg_min)
+                    bool const arg_min,
+                    rmm::cuda_stream_view stream)
     : num_rows(table.num_rows()),
-      comp(nullate::DYNAMIC{has_nulls}, table, table, nullptr, null_precedence),
+      comp(nullate::DYNAMIC{has_nulls}, table, table, nullptr, null_precedence, stream),
       arg_min(arg_min)
   {
   }
@@ -117,9 +118,10 @@ class comparison_binop_generator {
   }
 
  public:
-  auto binop() const
+  auto binop(rmm::cuda_stream_view stream) const
   {
-    return row_arg_minmax_fn(*d_flattened_input_ptr, has_nulls, null_orders_dvec.data(), is_min_op);
+    return row_arg_minmax_fn(
+      *d_flattened_input_ptr, has_nulls, null_orders_dvec.data(), is_min_op, stream);
   }
 
   template <typename BinOp>

cpp/src/search/search.cu

@@ -128,7 +128,8 @@ std::unique_ptr<column> search_ordered(table_view const& t,
                                                  lhs,
                                                  rhs,
                                                  column_order_dv.data(),
-                                                 null_precedence_dv.data());
+                                                 null_precedence_dv.data(),
+                                                 stream);
   launch_search(
     count_it, count_it, t.num_rows(), values.num_rows(), result_out, comp, find_first, stream);
 

cpp/src/sort/is_sorted.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -51,7 +51,8 @@ auto is_sorted(cudf::table_view const& in,
                                                  *d_input,
                                                  *d_input,
                                                  d_column_order.data(),
-                                                 d_null_precedence.data());
+                                                 d_null_precedence.data(),
+                                                 stream);
 
   auto sorted = thrust::is_sorted(rmm::exec_policy(stream),
                                   thrust::make_counting_iterator(0),

cpp/src/sort/sort_impl.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -132,7 +132,8 @@ std::unique_ptr<column> sorted_order(table_view input,
                                                        *device_table,
                                                        *device_table,
                                                        d_column_order.data(),
-                                                       d_null_precedence.data());
+                                                       d_null_precedence.data(),
+                                                       stream);
   if (stable) {
     thrust::stable_sort(rmm::exec_policy(stream),
                         mutable_indices_view.begin<size_type>(),

cpp/src/table/table_device_view.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -20,9 +20,12 @@
 #include <cudf/utilities/error.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/logical.h>
 
+#include <algorithm>
+
 namespace cudf {
 namespace detail {
 template <typename ColumnDeviceView, typename HostTableView>
@@ -67,15 +70,13 @@ struct is_relationally_comparable_functor {
 template <typename TableView>
 bool is_relationally_comparable(TableView const& lhs, TableView const& rhs)
 {
-  return thrust::all_of(thrust::counting_iterator<size_type>(0),
-                        thrust::counting_iterator<size_type>(lhs.num_columns()),
-                        [lhs, rhs] __device__(auto const i) {
-                          // Simplified this for compile time. (Ideally use double_type_dispatcher)
-                          // TODO: possible to implement without double type dispatcher.
-                          return lhs.column(i).type() == rhs.column(i).type() and
-                                 type_dispatcher(lhs.column(i).type(),
-                                                 is_relationally_comparable_functor{});
-                        });
+  return std::all_of(thrust::counting_iterator<size_type>(0),
+                     thrust::counting_iterator<size_type>(lhs.num_columns()),
+                     [lhs, rhs](auto const i) {
+                       return lhs.column(i).type() == rhs.column(i).type() and
+                              type_dispatcher(lhs.column(i).type(),
+                                              is_relationally_comparable_functor{});
+                     });
 }
 
 // Explicit extern template instantiation for a table of immutable views
@@ -86,13 +87,31 @@ extern template bool is_relationally_comparable<table_view>(table_view const& lh
 extern template bool is_relationally_comparable<mutable_table_view>(mutable_table_view const& lhs,
                                                                     mutable_table_view const& rhs);
 
+template <typename TableDeviceView>
+bool is_relationally_comparable(TableDeviceView const& lhs,
+                                TableDeviceView const& rhs,
+                                rmm::cuda_stream_view stream)
+{
+  return thrust::all_of(rmm::exec_policy(stream),
+                        thrust::counting_iterator<size_type>(0),
+                        thrust::counting_iterator<size_type>(lhs.num_columns()),
+                        [lhs, rhs] __device__(auto const i) {
+                          return lhs.column(i).type() == rhs.column(i).type() and
+                                 type_dispatcher(lhs.column(i).type(),
+                                                 is_relationally_comparable_functor{});
+                        });
+}
+
 // Explicit extern template instantiation for a device table of immutable views
 template bool is_relationally_comparable<table_device_view>(table_device_view const& lhs,
-                                                            table_device_view const& rhs);
+                                                            table_device_view const& rhs,
+                                                            rmm::cuda_stream_view stream);
 
 // Explicit extern template instantiation for a device table of mutable views
 template bool is_relationally_comparable<mutable_table_device_view>(
-  mutable_table_device_view const& lhs, mutable_table_device_view const& rhs);
+  mutable_table_device_view const& lhs,
+  mutable_table_device_view const& rhs,
+  rmm::cuda_stream_view stream);
 
 }  // namespace detail
 }  // namespace cudf