Explicitly disable groupby on unsupported key types.

cpp/include/cudf/table/row_operators.cuh

@@ -25,6 +25,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <thrust/equal.h>
+#include <thrust/logical.h>
 #include <thrust/swap.h>
 #include <thrust/transform_reduce.h>
 
@@ -222,13 +223,75 @@ class element_equality_comparator {
   bool nulls_are_equal;
 };
 
+namespace detail {
+
+/**
+ * @brief Helper functor to check if two types `LhsType` and `RhsType` are comparable for equality.
+ */
+template <typename LhsType>
+struct equality_comparable_functor {
+  template <typename RhsType>
+  bool __device__ operator()() const
+  {
+    return cudf::is_equality_comparable<LhsType, RhsType>();
+  }
+};
+
+/**
+ * @brief Helper functor to check if a specified column `rhs_col` can be compared for equality
+ * with type `LhsType`.
+ */
+struct lhs_equality_comparable_functor {
+  template <typename LhsType>
+  bool __device__ operator()(column_device_view rhs_col) const
+  {
+    return cudf::type_dispatcher(rhs_col.type(), equality_comparable_functor<LhsType>{});
+  }
+};
+
+/**
+ * @brief Checks if two columns have types that permit equality comparisons
+ *
+ * @param lhs Left column for equality comparisons
+ * @param rhs Right column for equality comparisons
+ * @return true If `lhs` and `rhs` columns have types whose values may be compared for equality
+ * @return false Otherwise.
+ */
+bool __device__ columns_are_equality_comparable(column_device_view lhs, column_device_view rhs)
+{
+  return cudf::type_dispatcher(lhs.type(), lhs_equality_comparable_functor{}, rhs);
+}
+
+/**
+ * @brief Checks if corresponding columns in `lhs` and `rhs` tables may be compared for equality
+ *
+ * @param lhs Left table, whose columns are to be compared for equality.
+ * @param rhs Right table, whose columns are to be compared for equality.
+ * @return true If *all* columns in `lhs` permit equality comparisons with corresponding columns in
+ * `rhs`.
+ * @return false Otherwise.
+ */
+bool all_columns_are_equality_comparable(table_device_view lhs, table_device_view rhs)
+{
+  auto columns_begin = thrust::make_counting_iterator<size_type>(0);
+  auto columns_end   = columns_begin + lhs.num_columns();
+  auto eq_comparable = [lhs, rhs] __device__(auto col_idx) {
+    return columns_are_equality_comparable(lhs.column(col_idx), rhs.column(col_idx));
+  };
+  return thrust::all_of(thrust::seq, columns_begin, columns_end, eq_comparable);
+}
+
+}  // namespace detail
+
 template <bool has_nulls = true>
 class row_equality_comparator {
  public:
   row_equality_comparator(table_device_view lhs, table_device_view rhs, bool nulls_are_equal = true)
     : lhs{lhs}, rhs{rhs}, nulls_are_equal{nulls_are_equal}
   {
     CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns.");
+    CUDF_EXPECTS(all_columns_are_equality_comparable(lhs, rhs),
+                 "Incompatible columns in row equality comparison.");
   }
 
   __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept

cpp/tests/CMakeLists.txt

@@ -62,6 +62,7 @@ ConfigureTest(GROUPBY_TEST
     groupby/count_tests.cpp
     groupby/groups_tests.cpp
     groupby/keys_tests.cpp
+    groupby/lists_tests.cpp
     groupby/m2_tests.cpp
     groupby/min_tests.cpp
     groupby/max_scan_tests.cpp

cpp/tests/groupby/lists_tests.cpp

@@ -0,0 +1,69 @@
+/*
+ * 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.
+ * 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 <tests/groupby/groupby_test_util.hpp>
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/iterator_utilities.hpp>
+#include <cudf_test/type_lists.hpp>
+
+#include <cudf/detail/aggregation/aggregation.hpp>
+
+namespace cudf {
+namespace test {
+
+template <typename V>
+struct groupby_lists_test : public cudf::test::BaseFixture {
+};
+
+TYPED_TEST_SUITE(groupby_lists_test, cudf::test::FixedWidthTypes);
+
+namespace {
+// Checking with a single aggregation, and aggregation column.
+// This test is orthogonal to the aggregation type; it focuses on testing the grouping
+// with LISTS keys.
+auto sum_agg() { return cudf::make_sum_aggregation<groupby_aggregation>(); }
+
+void test_sort_based_sum_agg(column_view const& keys, column_view const& values)
+{
+  test_single_agg(
+    keys, values, keys, values, sum_agg(), force_use_sort_impl::YES, null_policy::INCLUDE);
+}
+
+void test_hash_based_sum_agg(column_view const& keys, column_view const& values)
+{
+  test_single_agg(
+    keys, values, keys, values, sum_agg(), force_use_sort_impl::NO, null_policy::INCLUDE);
+}
+
+}  // namespace
+
+TYPED_TEST(groupby_lists_test, top_level_lists_are_unsupported)
+{
+  // Test that grouping on LISTS columns fails visibly.
+
+  // clang-format off
+  auto keys   = lists_column_wrapper<TypeParam, int32_t> { {1,1},  {2,2},  {3,3},   {1,1},   {2,2} };
+  auto values = fixed_width_column_wrapper<int32_t>      {     0,      1,      2,      3,       4  };
+  // clang-format on
+
+  EXPECT_THROW(test_sort_based_sum_agg(keys, values), cudf::logic_error);
+  EXPECT_THROW(test_hash_based_sum_agg(keys, values), cudf::logic_error);
+}
+
+}  // namespace test
+}  // namespace cudf

cpp/tests/groupby/structs_tests.cpp

@@ -22,8 +22,6 @@
 #include <cudf_test/type_lists.hpp>
 
 #include <cudf/detail/aggregation/aggregation.hpp>
-#include "cudf/aggregation.hpp"
-#include "cudf/types.hpp"
 
 using namespace cudf::test::iterators;
 
@@ -34,7 +32,7 @@ template <typename V>
 struct groupby_structs_test : public cudf::test::BaseFixture {
 };
 
-TYPED_TEST_CASE(groupby_structs_test, cudf::test::FixedWidthTypes);
+TYPED_TEST_SUITE(groupby_structs_test, cudf::test::FixedWidthTypes);
 
 using V       = int32_t;  // Type of Aggregation Column.
 using M0      = int32_t;  // Type of STRUCT's first (i.e. 0th) member.
@@ -79,27 +77,43 @@ void print_agg_results(column_view const& keys, column_view const& vals)
   }
 }
 
-void test_sum_agg(column_view const& keys,
-                  column_view const& values,
-                  column_view const& expected_keys,
-                  column_view const& expected_values)
+void test_sort_based_sum_agg(column_view const& keys,
+                             column_view const& values,
+                             column_view const& expected_keys,
+                             column_view const& expected_values)
 {
   test_single_agg(keys,
                   values,
                   expected_keys,
                   expected_values,
                   sum_agg(),
-                  force_use_sort_impl::NO,
+                  force_use_sort_impl::YES,
                   null_policy::INCLUDE);
+}
+
+void test_hash_based_sum_agg(column_view const& keys,
+                             column_view const& values,
+                             column_view const& expected_keys,
+                             column_view const& expected_values)
+{
   test_single_agg(keys,
                   values,
                   expected_keys,
                   expected_values,
                   sum_agg(),
-                  force_use_sort_impl::YES,
+                  force_use_sort_impl::NO,
                   null_policy::INCLUDE);
 }
 
+void test_sum_agg(column_view const& keys,
+                  column_view const& values,
+                  column_view const& expected_keys,
+                  column_view const& expected_values)
+{
+  test_sort_based_sum_agg(keys, values, expected_keys, expected_values);
+  test_hash_based_sum_agg(keys, values, expected_keys, expected_values);
+}
+
 }  // namespace
 
 TYPED_TEST(groupby_structs_test, basic)
@@ -312,7 +326,8 @@ TYPED_TEST(groupby_structs_test, lists_are_unsupported)
   // clang-format on
   auto keys = structs{{member_0, member_1}};
 
-  EXPECT_THROW(test_sum_agg(keys, values, keys, values), cudf::logic_error);
+  EXPECT_THROW(test_sort_based_sum_agg(keys, values, keys, values), cudf::logic_error);
+  EXPECT_THROW(test_hash_based_sum_agg(keys, values, keys, values), cudf::logic_error);
 }
 
 }  // namespace test