From 60b1c11fbd4b54a908969d570cd5a8ec03567a73 Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Tue, 28 Apr 2026 16:35:06 +0100
Subject: [PATCH 01/10] Implement equality of two table_views

At various times, it is useful to check whether two tables are equal. For
example, in cudf-polars we use this to check if two tables are
"compatibly" partitioned. Previously there have been no such utilities in
libcudf proper. The best one can do is to loop over the columns, call
cudf::binary_operation with NULL_EQUALS and then cudf::reduce on the
result. This launches many more kernels than necessary.

Instead, use the existing row_equality operators to perform a single
transform_reduce over the table checking for equality.
---
 cpp/CMakeLists.txt                  |   1 +
 cpp/include/cudf/table/equality.hpp |  35 ++++++++++
 cpp/src/table/table_equal.cu        |  75 ++++++++++++++++++++
 cpp/tests/table/table_tests.cpp     | 104 +++++++++++++++++++++++++++-
 4 files changed, 213 insertions(+), 2 deletions(-)
 create mode 100644 cpp/include/cudf/table/equality.hpp
 create mode 100644 cpp/src/table/table_equal.cu

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 5a0b2f95e830..91e8e86683c9 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -813,6 +813,7 @@ add_library(
   src/structs/utilities.cu
   src/table/table.cpp
   src/table/table_device_view.cu
+  src/table/table_equal.cu
   src/table/table_view.cpp
   src/text/deduplicate.cu
   src/text/detokenize.cu
diff --git a/cpp/include/cudf/table/equality.hpp b/cpp/include/cudf/table/equality.hpp
new file mode 100644
index 000000000000..83b442eca3ee
--- /dev/null
+++ b/cpp/include/cudf/table/equality.hpp
@@ -0,0 +1,35 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/table/table_view.hpp>
+#include <cudf/types.hpp>
+#include <cudf/utilities/default_stream.hpp>
+#include <cudf/utilities/memory_resource.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+namespace CUDF_EXPORT cudf {
+
+/**
+ * @brief Check if two tables are equal.
+ *
+ * Returns true if the input tables have the same number of rows, the same number of columns,
+ * matching column types, and every row in `left` compares equal to the row at the same index in
+ * `right`. Null equality is controlled by `nulls_equal`. Floating point NaN values compare equal.
+ *
+ * @throws cudf::logic_error if the tables contain non-equality-comparable column types.
+ *
+ * @param left The first table to compare
+ * @param right The second table to compare
+ * @param nulls_equal Flag to denote if null elements should be considered equal
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @return true if the tables are equal, false otherwise
+ */
+bool tables_equal(table_view const& left,
+                  table_view const& right,
+                  null_equality nulls_equal    = null_equality::EQUAL,
+                  rmm::cuda_stream_view stream = cudf::get_default_stream());
+
+}  // namespace CUDF_EXPORT cudf
diff --git a/cpp/src/table/table_equal.cu b/cpp/src/table/table_equal.cu
new file mode 100644
index 000000000000..0e340d202fe0
--- /dev/null
+++ b/cpp/src/table/table_equal.cu
@@ -0,0 +1,75 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/row_operator/common_utils.cuh>
+#include <cudf/detail/row_operator/equality.cuh>
+#include <cudf/table/equality.hpp>
+#include <cudf/table/table_view.hpp>
+#include <cudf/utilities/error.hpp>
+#include <cudf/utilities/type_checks.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <cuda/iterator>
+#include <cuda/std/functional>
+#include <thrust/transform_reduce.h>
+
+namespace cudf {
+namespace detail {
+namespace {
+
+template <bool has_nested_columns>
+bool tables_equal(table_view const& left,
+                  table_view const& right,
+                  null_equality nulls_equal,
+                  rmm::cuda_stream_view stream)
+{
+  auto const comparator = detail::row::equality::two_table_comparator{left, right, stream};
+  auto const rows_equal = comparator.equal_to<has_nested_columns>(
+    nullate::DYNAMIC{has_nested_nulls(left) or has_nested_nulls(right)}, nulls_equal);
+
+  return thrust::transform_reduce(
+    rmm::exec_policy_nosync(stream),
+    cuda::counting_iterator<size_type>{0},
+    cuda::counting_iterator<size_type>{left.num_rows()},
+    [rows_equal] __device__(size_type i) -> bool {
+      return rows_equal(detail::row::lhs_index_type{i}, detail::row::rhs_index_type{i});
+    },
+    true,
+    cuda::std::logical_and<bool>{});
+}
+
+}  // namespace
+
+bool tables_equal(table_view const& left,
+                  table_view const& right,
+                  null_equality nulls_equal,
+                  rmm::cuda_stream_view stream)
+{
+  if (left.num_rows() != right.num_rows() || left.num_columns() != right.num_columns() ||
+      !have_same_types(left, right)) {
+    return false;
+  } else if (left.num_rows() == 0) {
+    return true;
+  }
+
+  return cudf::has_nested_columns(left) || cudf::has_nested_columns(right)
+           ? tables_equal<true>(left, right, nulls_equal, stream)
+           : tables_equal<false>(left, right, nulls_equal, stream);
+}
+}  // namespace detail
+
+bool tables_equal(table_view const& left,
+                  table_view const& right,
+                  null_equality nulls_equal,
+                  rmm::cuda_stream_view stream)
+{
+  CUDF_FUNC_RANGE();
+  return detail::tables_equal(left, right, nulls_equal, stream);
+}
+
+}  // namespace cudf
diff --git a/cpp/tests/table/table_tests.cpp b/cpp/tests/table/table_tests.cpp
index b909d9c9392f..725f90af1623 100644
--- a/cpp/tests/table/table_tests.cpp
+++ b/cpp/tests/table/table_tests.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -10,15 +10,20 @@
 
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_view.hpp>
+#include <cudf/table/equality.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/table/table_view.hpp>
 
+#include <rmm/device_buffer.hpp>
+
+#include <limits>
 #include <memory>
 
 template <typename T>
 using column_wrapper = cudf::test::fixed_width_column_wrapper<T>;
 
-using s_col_wrapper = cudf::test::strings_column_wrapper;
+using s_col_wrapper       = cudf::test::strings_column_wrapper;
+using structs_col_wrapper = cudf::test::structs_column_wrapper;
 
 using CVector     = std::vector<std::unique_ptr<cudf::column>>;
 using column      = cudf::column;
@@ -163,4 +168,99 @@ TEST_F(TableTest, AllocSizeWithNulls)
   EXPECT_EQ(t.alloc_size(), 152);  // bitmask has padding
 }
 
+TEST_F(TableTest, TablesEqual)
+{
+  column_wrapper<int32_t> left_col0{{1, 2, 3}};
+  column_wrapper<double> left_col1{{4.0, 5.0, 6.0}};
+  column_wrapper<int32_t> right_col0{{1, 2, 3}};
+  column_wrapper<double> right_col1{{4.0, 5.0, 6.0}};
+
+  EXPECT_TRUE(cudf::tables_equal(cudf::table_view{{left_col0, left_col1}},
+                                 cudf::table_view{{right_col0, right_col1}}));
+}
+
+TEST_F(TableTest, TablesEqualValueMismatch)
+{
+  column_wrapper<int32_t> left{{1, 2, 3}};
+  column_wrapper<int32_t> right{{1, 4, 3}};
+
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}));
+}
+
+TEST_F(TableTest, TablesEqualShapeAndTypeMismatch)
+{
+  column_wrapper<int32_t> left{{1, 2, 3}};
+  column_wrapper<int32_t> shorter{{1, 2}};
+  column_wrapper<int32_t> extra{{1, 2, 3}};
+  column_wrapper<int64_t> different_type{{1, 2, 3}};
+
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{shorter}}));
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{left, extra}}));
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{different_type}}));
+}
+
+TEST_F(TableTest, TablesEqualNullEquality)
+{
+  column_wrapper<int32_t> left{{1, 2, 3}, {1, 0, 1}};
+  column_wrapper<int32_t> right{{1, 4, 3}, {1, 0, 1}};
+
+  EXPECT_TRUE(cudf::tables_equal(
+    cudf::table_view{{left}}, cudf::table_view{{right}}, cudf::null_equality::EQUAL));
+  EXPECT_FALSE(cudf::tables_equal(
+    cudf::table_view{{left}}, cudf::table_view{{right}}, cudf::null_equality::UNEQUAL));
+}
+
+TEST_F(TableTest, TablesEqualNaNsCompareEqual)
+{
+  column_wrapper<double> left{{std::numeric_limits<double>::quiet_NaN(), 1.0}};
+  column_wrapper<double> right{{std::numeric_limits<double>::quiet_NaN(), 1.0}};
+
+  EXPECT_TRUE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}));
+}
+
+TEST_F(TableTest, TablesEqualStructColumns)
+{
+  column_wrapper<int32_t> left_id{{1, 2, 3}};
+  column_wrapper<int16_t> left_inner_value{{10, 20, 30}};
+  column_wrapper<double> left_deep_leaf{{1.25, 2.5, 3.75}};
+  structs_col_wrapper left_inner{{left_inner_value, left_deep_leaf}};
+  structs_col_wrapper left_outer{{left_id, left_inner}};
+
+  column_wrapper<int32_t> right_id{{1, 2, 3}};
+  column_wrapper<int16_t> right_inner_value{{10, 20, 30}};
+  column_wrapper<double> right_deep_leaf{{1.25, 2.5, 3.75}};
+  structs_col_wrapper right_inner{{right_inner_value, right_deep_leaf}};
+  structs_col_wrapper right_outer{{right_id, right_inner}};
+
+  EXPECT_TRUE(cudf::tables_equal(cudf::table_view{{left_outer}}, cudf::table_view{{right_outer}}));
+}
+
+TEST_F(TableTest, TablesEqualStructColumnsDeepLeafMismatch)
+{
+  column_wrapper<int32_t> left_id{{1, 2, 3}};
+  column_wrapper<int16_t> left_inner_value{{10, 20, 30}};
+  column_wrapper<double> left_deep_leaf{{1.25, 2.5, 3.75}};
+  structs_col_wrapper left_inner{{left_inner_value, left_deep_leaf}};
+  structs_col_wrapper left_outer{{left_id, left_inner}};
+
+  column_wrapper<int32_t> right_id{{1, 2, 3}};
+  column_wrapper<int16_t> right_inner_value{{10, 20, 30}};
+  column_wrapper<double> right_deep_leaf{{1.25, 2.5, 99.0}};
+  structs_col_wrapper right_inner{{right_inner_value, right_deep_leaf}};
+  structs_col_wrapper right_outer{{right_id, right_inner}};
+
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left_outer}}, cudf::table_view{{right_outer}}));
+}
+
+TEST_F(TableTest, TablesEqualThrowsForNonEqualityComparableTypes)
+{
+  auto left =
+    column{cudf::data_type{cudf::type_id::EMPTY}, 3, rmm::device_buffer{}, rmm::device_buffer{}, 0};
+  auto right =
+    column{cudf::data_type{cudf::type_id::EMPTY}, 3, rmm::device_buffer{}, rmm::device_buffer{}, 0};
+
+  EXPECT_THROW(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}),
+               cudf::logic_error);
+}
+
 CUDF_TEST_PROGRAM_MAIN()

From e23ea20156317b5381f48815534da59c94c2bd2f Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Wed, 29 Apr 2026 10:09:39 +0100
Subject: [PATCH 02/10] Use transform then reduce

The row_operator function is too complex for transform-reduce, resulting in
very long compile times in general, and a bug in cicc 13.1.
---
 cpp/src/table/table_equal.cu | 19 ++++++++++++-------
 1 file changed, 12 insertions(+), 7 deletions(-)

diff --git a/cpp/src/table/table_equal.cu b/cpp/src/table/table_equal.cu
index 0e340d202fe0..91b181a15f56 100644
--- a/cpp/src/table/table_equal.cu
+++ b/cpp/src/table/table_equal.cu
@@ -3,20 +3,23 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
+#include <cudf/detail/algorithms/reduce.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/row_operator/common_utils.cuh>
 #include <cudf/detail/row_operator/equality.cuh>
 #include <cudf/table/equality.hpp>
 #include <cudf/table/table_view.hpp>
 #include <cudf/utilities/error.hpp>
+#include <cudf/utilities/memory_resource.hpp>
 #include <cudf/utilities/type_checks.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cub/device/device_transform.cuh>
 #include <cuda/iterator>
 #include <cuda/std/functional>
-#include <thrust/transform_reduce.h>
 
 namespace cudf {
 namespace detail {
@@ -31,16 +34,18 @@ bool tables_equal(table_view const& left,
   auto const comparator = detail::row::equality::two_table_comparator{left, right, stream};
   auto const rows_equal = comparator.equal_to<has_nested_columns>(
     nullate::DYNAMIC{has_nested_nulls(left) or has_nested_nulls(right)}, nulls_equal);
-
-  return thrust::transform_reduce(
-    rmm::exec_policy_nosync(stream),
+  rmm::device_uvector<bool> eq_rows{
+    static_cast<std::size_t>(left.num_rows()), stream, cudf::get_current_device_resource_ref()};
+  CUDF_CUDA_TRY(cub::DeviceTransform::Transform(
     cuda::counting_iterator<size_type>{0},
-    cuda::counting_iterator<size_type>{left.num_rows()},
+    eq_rows.begin(),
+    eq_rows.size(),
     [rows_equal] __device__(size_type i) -> bool {
       return rows_equal(detail::row::lhs_index_type{i}, detail::row::rhs_index_type{i});
     },
-    true,
-    cuda::std::logical_and<bool>{});
+    stream.value()));
+  return cudf::detail::reduce(
+    eq_rows.begin(), eq_rows.end(), true, cuda::std::logical_and<bool>{}, stream);
 }
 
 }  // namespace

From 8141248e304c81e86b8f76b32428e9feb6995ee2 Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Wed, 29 Apr 2026 10:14:32 +0100
Subject: [PATCH 03/10] Throw documented exception in two_table_comparator

Previously if two tables had column types that were not
equality-comparable, cudf::logic_error was thrown, while the documented
exception was std::invalid_argument. Fix this by throwing the correct
exception.
---
 cpp/include/cudf/table/equality.hpp   |  2 +-
 cpp/src/row_operator/row_operators.cu | 10 ++++++----
 cpp/tests/table/table_tests.cpp       |  3 ++-
 3 files changed, 9 insertions(+), 6 deletions(-)

diff --git a/cpp/include/cudf/table/equality.hpp b/cpp/include/cudf/table/equality.hpp
index 83b442eca3ee..f1e806d82029 100644
--- a/cpp/include/cudf/table/equality.hpp
+++ b/cpp/include/cudf/table/equality.hpp
@@ -19,7 +19,7 @@ namespace CUDF_EXPORT cudf {
  * matching column types, and every row in `left` compares equal to the row at the same index in
  * `right`. Null equality is controlled by `nulls_equal`. Floating point NaN values compare equal.
  *
- * @throws cudf::logic_error if the tables contain non-equality-comparable column types.
+ * @throws cudf::logic_error if the tables contain `EMPTY` types.
  *
  * @param left The first table to compare
  * @param right The second table to compare
diff --git a/cpp/src/row_operator/row_operators.cu b/cpp/src/row_operator/row_operators.cu
index 9ab21e9cfcbf..4b21d9c980cb 100644
--- a/cpp/src/row_operator/row_operators.cu
+++ b/cpp/src/row_operator/row_operators.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -25,6 +25,7 @@
 #include <thrust/iterator/transform_iterator.h>
 
 #include <functional>
+#include <stdexcept>
 
 namespace cudf {
 namespace detail {
@@ -356,9 +357,10 @@ 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()));
+      CUDF_EXPECTS(
+        is_equality_comparable(c.type()),
+        "Cannot compare equality for a table with a column of type " + cudf::type_to_name(c.type()),
+        std::invalid_argument);
     }
     for (auto child = c.child_begin(); child < c.child_end(); ++child) {
       check_column(*child);
diff --git a/cpp/tests/table/table_tests.cpp b/cpp/tests/table/table_tests.cpp
index 725f90af1623..2bed02e45555 100644
--- a/cpp/tests/table/table_tests.cpp
+++ b/cpp/tests/table/table_tests.cpp
@@ -18,6 +18,7 @@
 
 #include <limits>
 #include <memory>
+#include <stdexcept>
 
 template <typename T>
 using column_wrapper = cudf::test::fixed_width_column_wrapper<T>;
@@ -260,7 +261,7 @@ TEST_F(TableTest, TablesEqualThrowsForNonEqualityComparableTypes)
     column{cudf::data_type{cudf::type_id::EMPTY}, 3, rmm::device_buffer{}, rmm::device_buffer{}, 0};
 
   EXPECT_THROW(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}),
-               cudf::logic_error);
+               std::invalid_argument);
 }
 
 CUDF_TEST_PROGRAM_MAIN()

From 4c875d0b5ed215d060bda5b2f8432f06ed166f8a Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Wed, 29 Apr 2026 11:02:09 +0100
Subject: [PATCH 04/10] More tests

---
 cpp/tests/table/table_tests.cpp | 33 ++++++++++++++++++++++++++++++++-
 1 file changed, 32 insertions(+), 1 deletion(-)

diff --git a/cpp/tests/table/table_tests.cpp b/cpp/tests/table/table_tests.cpp
index 2bed02e45555..40cf43ecce01 100644
--- a/cpp/tests/table/table_tests.cpp
+++ b/cpp/tests/table/table_tests.cpp
@@ -24,6 +24,7 @@ template <typename T>
 using column_wrapper = cudf::test::fixed_width_column_wrapper<T>;
 
 using s_col_wrapper       = cudf::test::strings_column_wrapper;
+using lists_col_wrapper   = cudf::test::lists_column_wrapper<int32_t>;
 using structs_col_wrapper = cudf::test::structs_column_wrapper;
 
 using CVector     = std::vector<std::unique_ptr<cudf::column>>;
@@ -261,7 +262,37 @@ TEST_F(TableTest, TablesEqualThrowsForNonEqualityComparableTypes)
     column{cudf::data_type{cudf::type_id::EMPTY}, 3, rmm::device_buffer{}, rmm::device_buffer{}, 0};
 
   EXPECT_THROW(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}),
-               std::invalid_argument);
+               cudf::logic_error);
+}
+
+TEST_F(TableTest, TablesEqualListColumns)
+{
+  lists_col_wrapper left{{1, 2}, {3}, {}};
+  lists_col_wrapper right{{1, 2}, {3}, {}};
+  lists_col_wrapper different_values{{1, 2}, {4}, {}};
+  lists_col_wrapper different_offsets{{1}, {2, 3}, {}};
+
+  EXPECT_TRUE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}));
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{different_values}}));
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{different_offsets}}));
+}
+
+TEST_F(TableTest, TablesEqualStructColumnsWithLists)
+{
+  column_wrapper<int32_t> left_id{{1, 2, 3}};
+  lists_col_wrapper left_list{{1, 2}, {3}, {}};
+  structs_col_wrapper left{{left_id, left_list}};
+
+  column_wrapper<int32_t> right_id{{1, 2, 3}};
+  lists_col_wrapper right_list{{1, 2}, {3}, {}};
+  structs_col_wrapper right{{right_id, right_list}};
+
+  column_wrapper<int32_t> different_id{{1, 2, 3}};
+  lists_col_wrapper different_list{{1, 2}, {4}, {}};
+  structs_col_wrapper different{{different_id, different_list}};
+
+  EXPECT_TRUE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}));
+  EXPECT_FALSE(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{different}}));
 }
 
 CUDF_TEST_PROGRAM_MAIN()

From a25682e0374bb415dd12d730e28738d6456df884 Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Wed, 29 Apr 2026 11:02:28 +0100
Subject: [PATCH 05/10] Fix traits test

---
 cpp/tests/types/traits_test.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/tests/types/traits_test.cpp b/cpp/tests/types/traits_test.cpp
index 6dbaadf56227..6ddbe7a53692 100644
--- a/cpp/tests/types/traits_test.cpp
+++ b/cpp/tests/types/traits_test.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -115,7 +115,7 @@ TYPED_TEST(TypedTraitsTest, NotEqualityComparableWithList)
   bool comparable = cudf::is_equality_comparable<TypeParam, cudf::list_view>();
   EXPECT_FALSE(comparable);
 
-  cudf::is_equality_comparable<cudf::list_view, cudf::list_view>();
+  comparable = cudf::is_equality_comparable<cudf::list_view, cudf::list_view>();
   EXPECT_FALSE(comparable);
 }
 

From 2d83e62797658442bf412a7e2e4a69963dd5eb82 Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Wed, 29 Apr 2026 14:34:33 +0100
Subject: [PATCH 06/10] Add streams test

---
 cpp/tests/CMakeLists.txt                  |  1 +
 cpp/tests/streams/table_equality_test.cpp | 28 +++++++++++++++++++++++
 2 files changed, 29 insertions(+)
 create mode 100644 cpp/tests/streams/table_equality_test.cpp

diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index b75b72b1aee3..3fc6dbf96739 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -790,6 +790,7 @@ ConfigureTest(
   STREAM_MODE
   testing
 )
+ConfigureTest(STREAM_TABLE_EQUALITY_TEST streams/table_equality_test.cpp STREAM_MODE testing)
 ConfigureTest(
   STREAM_TEXT_TEST
   streams/text/edit_distance_test.cpp
diff --git a/cpp/tests/streams/table_equality_test.cpp b/cpp/tests/streams/table_equality_test.cpp
new file mode 100644
index 000000000000..8c7879c4ec7e
--- /dev/null
+++ b/cpp/tests/streams/table_equality_test.cpp
@@ -0,0 +1,28 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/default_stream.hpp>
+#include <cudf_test/testing_main.hpp>
+
+#include <cudf/table/equality.hpp>
+#include <cudf/table/table_view.hpp>
+#include <cudf/types.hpp>
+
+class TableEqualTest : public cudf::test::BaseFixture {};
+
+TEST_F(TableEqualTest, NotEqual)
+{
+  cudf::test::fixed_width_column_wrapper<int> left(
+    {{0, 0, 0, 0, 0}, {false, false, true, true, true}});
+  cudf::test::fixed_width_column_wrapper<int> right({1, 1, 1, 1, 1});
+  cudf::tables_equal(cudf::table_view{{left}},
+                     cudf::table_view{{right}},
+                     cudf::null_equality::EQUAL,
+                     cudf::test::get_default_stream());
+}
+
+CUDF_TEST_PROGRAM_MAIN()

From d6126658e818cc88abff81351b3fef6c3d192bb3 Mon Sep 17 00:00:00 2001
From: Lawrence Mitchell <lmitchell@nvidia.com>
Date: Fri, 8 May 2026 10:52:28 +0100
Subject: [PATCH 07/10] Nodiscard

---
 cpp/include/cudf/table/equality.hpp       |  8 ++++----
 cpp/src/table/table_equal.cu              | 16 ++++++++--------
 cpp/tests/streams/table_equality_test.cpp |  8 ++++----
 cpp/tests/table/table_tests.cpp           |  5 +++--
 4 files changed, 19 insertions(+), 18 deletions(-)

diff --git a/cpp/include/cudf/table/equality.hpp b/cpp/include/cudf/table/equality.hpp
index f1e806d82029..22af77300f77 100644
--- a/cpp/include/cudf/table/equality.hpp
+++ b/cpp/include/cudf/table/equality.hpp
@@ -27,9 +27,9 @@ namespace CUDF_EXPORT cudf {
  * @param stream CUDA stream used for device memory operations and kernel launches
  * @return true if the tables are equal, false otherwise
  */
-bool tables_equal(table_view const& left,
-                  table_view const& right,
-                  null_equality nulls_equal    = null_equality::EQUAL,
-                  rmm::cuda_stream_view stream = cudf::get_default_stream());
+[[nodiscard]] bool tables_equal(table_view const& left,
+                                table_view const& right,
+                                null_equality nulls_equal    = null_equality::EQUAL,
+                                rmm::cuda_stream_view stream = cudf::get_default_stream());
 
 }  // namespace CUDF_EXPORT cudf
diff --git a/cpp/src/table/table_equal.cu b/cpp/src/table/table_equal.cu
index 91b181a15f56..0cc97e0da260 100644
--- a/cpp/src/table/table_equal.cu
+++ b/cpp/src/table/table_equal.cu
@@ -26,10 +26,10 @@ namespace detail {
 namespace {
 
 template <bool has_nested_columns>
-bool tables_equal(table_view const& left,
-                  table_view const& right,
-                  null_equality nulls_equal,
-                  rmm::cuda_stream_view stream)
+[[nodiscard]] bool tables_equal(table_view const& left,
+                                table_view const& right,
+                                null_equality nulls_equal,
+                                rmm::cuda_stream_view stream)
 {
   auto const comparator = detail::row::equality::two_table_comparator{left, right, stream};
   auto const rows_equal = comparator.equal_to<has_nested_columns>(
@@ -50,10 +50,10 @@ bool tables_equal(table_view const& left,
 
 }  // namespace
 
-bool tables_equal(table_view const& left,
-                  table_view const& right,
-                  null_equality nulls_equal,
-                  rmm::cuda_stream_view stream)
+[[nodiscard]] bool tables_equal(table_view const& left,
+                                table_view const& right,
+                                null_equality nulls_equal,
+                                rmm::cuda_stream_view stream)
 {
   if (left.num_rows() != right.num_rows() || left.num_columns() != right.num_columns() ||
       !have_same_types(left, right)) {
diff --git a/cpp/tests/streams/table_equality_test.cpp b/cpp/tests/streams/table_equality_test.cpp
index 8c7879c4ec7e..fdd85eb1b1b5 100644
--- a/cpp/tests/streams/table_equality_test.cpp
+++ b/cpp/tests/streams/table_equality_test.cpp
@@ -19,10 +19,10 @@ TEST_F(TableEqualTest, NotEqual)
   cudf::test::fixed_width_column_wrapper<int> left(
     {{0, 0, 0, 0, 0}, {false, false, true, true, true}});
   cudf::test::fixed_width_column_wrapper<int> right({1, 1, 1, 1, 1});
-  cudf::tables_equal(cudf::table_view{{left}},
-                     cudf::table_view{{right}},
-                     cudf::null_equality::EQUAL,
-                     cudf::test::get_default_stream());
+  std::ignore = cudf::tables_equal(cudf::table_view{{left}},
+                                   cudf::table_view{{right}},
+                                   cudf::null_equality::EQUAL,
+                                   cudf::test::get_default_stream());
 }
 
 CUDF_TEST_PROGRAM_MAIN()
diff --git a/cpp/tests/table/table_tests.cpp b/cpp/tests/table/table_tests.cpp
index 40cf43ecce01..c7d37b943b65 100644
--- a/cpp/tests/table/table_tests.cpp
+++ b/cpp/tests/table/table_tests.cpp
@@ -261,8 +261,9 @@ TEST_F(TableTest, TablesEqualThrowsForNonEqualityComparableTypes)
   auto right =
     column{cudf::data_type{cudf::type_id::EMPTY}, 3, rmm::device_buffer{}, rmm::device_buffer{}, 0};
 
-  EXPECT_THROW(cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}),
-               cudf::logic_error);
+  EXPECT_THROW(
+    std::ignore = cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}),
+    cudf::logic_error);
 }
 
 TEST_F(TableTest, TablesEqualListColumns)

From 13cdd7798189b65a5bfaf39d78c84ed2e87f5412 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <mhaseeb@nvidia.com>
Date: Fri, 8 May 2026 19:01:43 +0000
Subject: [PATCH 08/10] Use table_equal API in examples

---
 .agents/skills/build-test-cudf/SKILL.md      |  9 ++++++
 .devcontainer/cuda13.1-pip/devcontainer.json |  5 ++--
 .gitignore                                   |  3 ++
 cpp/examples/hybrid_scan_io/common_utils.cpp | 31 ++++----------------
 cpp/examples/parquet_io/common_utils.cpp     | 19 ++++--------
 cpp/examples/parquet_io/parquet_io.cpp       |  1 -
 6 files changed, 25 insertions(+), 43 deletions(-)

diff --git a/.agents/skills/build-test-cudf/SKILL.md b/.agents/skills/build-test-cudf/SKILL.md
index d10738b0a0dd..5f722d429d22 100644
--- a/.agents/skills/build-test-cudf/SKILL.md
+++ b/.agents/skills/build-test-cudf/SKILL.md
@@ -18,6 +18,15 @@ cudf packages must be built in this order as needed (each depends on the previou
 
 When rebuilding a package, ensure its dependencies are already built. Since we are in cudf devcontainer, DO NOT run the `build.sh` script or install any packages yourself. Instead, always use the `build-*`, `test-*`, `rapids-*`, `clean-*` scripts located at `/usr/bin`.
 
+### Set native architecture
+
+Set and export `CUDAARCHS=NATIVE` before building cudf so we only build it for the GPU architectures installed on the machine.
+
+```bash
+export CUDAARCHS=NATIVE # Direct export
+set-native              # Alternative: alias defined in .bashrc
+```
+
 ### Building libcudf
 
 Whenever building libcudf from scratch (CMake has not been run before), run:
diff --git a/.devcontainer/cuda13.1-pip/devcontainer.json b/.devcontainer/cuda13.1-pip/devcontainer.json
index 730b1c1e8cae..70bb71f52110 100644
--- a/.devcontainer/cuda13.1-pip/devcontainer.json
+++ b/.devcontainer/cuda13.1-pip/devcontainer.json
@@ -11,7 +11,7 @@
   "runArgs": [
     "--rm",
     "--name",
-    "${localEnv:USER:anon}-rapids-${localWorkspaceFolderBasename}-26.06-cuda13.1-pip",
+    "${localEnv:USER:anon}-starburst-${localWorkspaceFolderBasename}-26.06-cuda13.1-pip",
     "--ulimit",
     "nofile=500000"
   ],
@@ -41,7 +41,8 @@
     "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/../.local/share/${localWorkspaceFolderBasename}-cuda13.1-venvs,target=/home/coder/.local/share/venvs,type=bind,consistency=consistent",
-    "source=${localWorkspaceFolder}/../rapidsmpf,target=/home/coder/rapidsmpf,type=bind,consistency=consistent"
+    "source=${localWorkspaceFolder}/../rapidsmpf,target=/home/coder/rapidsmpf,type=bind,consistency=consistent",
+    "source=/raid/mhaseeb/data,target=/data,type=bind,readonly"
   ],
   "customizations": {
     "vscode": {
diff --git a/.gitignore b/.gitignore
index 180a6a286e2a..f3959b83a954 100644
--- a/.gitignore
+++ b/.gitignore
@@ -181,3 +181,6 @@ compile_commands.json
 # pytest artifacts
 rmm_log.txt
 python/cudf/cudf_pandas_tests/data/rmm_log.txt
+
+# agents
+.agents/
diff --git a/cpp/examples/hybrid_scan_io/common_utils.cpp b/cpp/examples/hybrid_scan_io/common_utils.cpp
index 635686171866..6f08f8e1965d 100644
--- a/cpp/examples/hybrid_scan_io/common_utils.cpp
+++ b/cpp/examples/hybrid_scan_io/common_utils.cpp
@@ -6,8 +6,7 @@
 #include "common_utils.hpp"
 
 #include <cudf/ast/expressions.hpp>
-#include <cudf/io/parquet.hpp>
-#include <cudf/join/filtered_join.hpp>
+#include <cudf/table/equality.hpp>
 #include <cudf/table/table_view.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -69,30 +68,10 @@ void check_tables_equal(cudf::table_view const& lhs_table,
                         cudf::table_view const& rhs_table,
                         rmm::cuda_stream_view stream)
 {
-  try {
-    // Left anti-join the original and transcoded tables identical tables should not throw an
-    // exception and return an empty indices vector
-    cudf::filtered_join join_obj(lhs_table, cudf::null_equality::EQUAL, stream);
-    auto const indices = join_obj.anti_join(rhs_table, stream);
-    // No exception thrown, check indices
-    auto const tables_equal = indices->size() == 0;
-    if (tables_equal) {
-      std::cout << "Tables identical: " << std::boolalpha << tables_equal << "\n\n";
-    } else {
-      // Helper to write parquet data for inspection
-      auto const write_parquet =
-        [](cudf::table_view table, std::string filepath, rmm::cuda_stream_view stream) {
-          auto sink_info = cudf::io::sink_info(filepath);
-          auto opts      = cudf::io::parquet_writer_options::builder(sink_info, table).build();
-          cudf::io::write_parquet(opts, stream);
-        };
-      write_parquet(lhs_table, "lhs_table.parquet", stream);
-      write_parquet(rhs_table, "rhs_table.parquet", stream);
-      throw std::logic_error("Tables identical: false\n\n");
-    }
-  } catch (std::exception& e) {
-    std::cout << e.what() << std::endl;
-  }
+  auto const tables_equal =
+    cudf::tables_equal(lhs_table, rhs_table, cudf::null_equality::EQUAL, stream);
+  std::cout << "Tables identical: " << std::boolalpha << tables_equal << "\n\n";
+  if (not tables_equal) { throw std::logic_error("Table equality check failed"); }
 }
 
 std::vector<io_source> extract_input_sources(std::string const& paths,
diff --git a/cpp/examples/parquet_io/common_utils.cpp b/cpp/examples/parquet_io/common_utils.cpp
index 99c2df7797fa..251391537ab2 100644
--- a/cpp/examples/parquet_io/common_utils.cpp
+++ b/cpp/examples/parquet_io/common_utils.cpp
@@ -7,7 +7,7 @@
 
 #include <cudf/concatenate.hpp>
 #include <cudf/io/types.hpp>
-#include <cudf/join/filtered_join.hpp>
+#include <cudf/table/equality.hpp>
 #include <cudf/table/table_view.hpp>
 
 #include <rmm/mr/cuda_memory_resource.hpp>
@@ -85,19 +85,10 @@ void check_tables_equal(cudf::table_view const& lhs_table,
                         cudf::table_view const& rhs_table,
                         rmm::cuda_stream_view stream)
 {
-  try {
-    // Left anti-join the original and transcoded tables identical tables should not throw an
-    // exception and return an empty indices vector
-    cudf::filtered_join join_obj(lhs_table, cudf::null_equality::EQUAL, stream);
-    auto const indices = join_obj.anti_join(rhs_table, stream);
-
-    // No exception thrown, check indices
-    auto const valid = indices->size() == 0;
-    std::cout << "Tables identical: " << valid << "\n\n";
-  } catch (std::exception& e) {
-    std::cerr << e.what() << std::endl << std::endl;
-    throw std::runtime_error("Tables identical: false\n\n");
-  }
+  auto const tables_equal =
+    cudf::tables_equal(lhs_table, rhs_table, cudf::null_equality::EQUAL, stream);
+  std::cout << "Tables identical: " << std::boolalpha << tables_equal << "\n\n";
+  if (not tables_equal) { throw std::logic_error("Table equality check failed"); }
 }
 
 std::unique_ptr<cudf::table> concatenate_tables(std::vector<std::unique_ptr<cudf::table>> tables,
diff --git a/cpp/examples/parquet_io/parquet_io.cpp b/cpp/examples/parquet_io/parquet_io.cpp
index a80ec5d44d54..33f78f060a03 100644
--- a/cpp/examples/parquet_io/parquet_io.cpp
+++ b/cpp/examples/parquet_io/parquet_io.cpp
@@ -4,7 +4,6 @@
  */
 
 #include "common_utils.hpp"
-#include "io_source.hpp"
 #include "timer.hpp"
 
 #include <cudf/io/parquet.hpp>

From 7eab00276aef9e3b99fc4349eacf2d065349f77e Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 8 May 2026 12:04:20 -0700
Subject: [PATCH 09/10] Update .gitignore

---
 .gitignore | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/.gitignore b/.gitignore
index f3959b83a954..180a6a286e2a 100644
--- a/.gitignore
+++ b/.gitignore
@@ -181,6 +181,3 @@ compile_commands.json
 # pytest artifacts
 rmm_log.txt
 python/cudf/cudf_pandas_tests/data/rmm_log.txt
-
-# agents
-.agents/

From 3b8fd3d2a877c27d59491e70b0f0913007e29fec Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <mhaseeb@nvidia.com>
Date: Fri, 8 May 2026 19:06:50 +0000
Subject: [PATCH 10/10] Undo erroneous changes

---
 .agents/skills/build-test-cudf/SKILL.md      | 9 ---------
 .devcontainer/cuda13.1-pip/devcontainer.json | 5 ++---
 2 files changed, 2 insertions(+), 12 deletions(-)

diff --git a/.agents/skills/build-test-cudf/SKILL.md b/.agents/skills/build-test-cudf/SKILL.md
index 5f722d429d22..d10738b0a0dd 100644
--- a/.agents/skills/build-test-cudf/SKILL.md
+++ b/.agents/skills/build-test-cudf/SKILL.md
@@ -18,15 +18,6 @@ cudf packages must be built in this order as needed (each depends on the previou
 
 When rebuilding a package, ensure its dependencies are already built. Since we are in cudf devcontainer, DO NOT run the `build.sh` script or install any packages yourself. Instead, always use the `build-*`, `test-*`, `rapids-*`, `clean-*` scripts located at `/usr/bin`.
 
-### Set native architecture
-
-Set and export `CUDAARCHS=NATIVE` before building cudf so we only build it for the GPU architectures installed on the machine.
-
-```bash
-export CUDAARCHS=NATIVE # Direct export
-set-native              # Alternative: alias defined in .bashrc
-```
-
 ### Building libcudf
 
 Whenever building libcudf from scratch (CMake has not been run before), run:
diff --git a/.devcontainer/cuda13.1-pip/devcontainer.json b/.devcontainer/cuda13.1-pip/devcontainer.json
index 70bb71f52110..730b1c1e8cae 100644
--- a/.devcontainer/cuda13.1-pip/devcontainer.json
+++ b/.devcontainer/cuda13.1-pip/devcontainer.json
@@ -11,7 +11,7 @@
   "runArgs": [
     "--rm",
     "--name",
-    "${localEnv:USER:anon}-starburst-${localWorkspaceFolderBasename}-26.06-cuda13.1-pip",
+    "${localEnv:USER:anon}-rapids-${localWorkspaceFolderBasename}-26.06-cuda13.1-pip",
     "--ulimit",
     "nofile=500000"
   ],
@@ -41,8 +41,7 @@
     "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent",
     "source=${localWorkspaceFolder}/../.local/share/${localWorkspaceFolderBasename}-cuda13.1-venvs,target=/home/coder/.local/share/venvs,type=bind,consistency=consistent",
-    "source=${localWorkspaceFolder}/../rapidsmpf,target=/home/coder/rapidsmpf,type=bind,consistency=consistent",
-    "source=/raid/mhaseeb/data,target=/data,type=bind,readonly"
+    "source=${localWorkspaceFolder}/../rapidsmpf,target=/home/coder/rapidsmpf,type=bind,consistency=consistent"
   ],
   "customizations": {
     "vscode": {