diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 891a9b202b90..030023801a5b 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -894,6 +894,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..22af77300f77
--- /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 `EMPTY` 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
+ */
+[[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/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/src/table/table_equal.cu b/cpp/src/table/table_equal.cu
new file mode 100644
index 000000000000..0cc97e0da260
--- /dev/null
+++ b/cpp/src/table/table_equal.cu
@@ -0,0 +1,80 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * 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>
+
+namespace cudf {
+namespace detail {
+namespace {
+
+template <bool has_nested_columns>
+[[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>(
+    nullate::DYNAMIC{has_nested_nulls(left) or has_nested_nulls(right)}, nulls_equal);
+  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},
+    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});
+    },
+    stream.value()));
+  return cudf::detail::reduce(
+    eq_rows.begin(), eq_rows.end(), true, cuda::std::logical_and<bool>{}, stream);
+}
+
+}  // namespace
+
+[[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)) {
+    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/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index 3f95595ad6f6..bae781f36bda 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -793,6 +793,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..fdd85eb1b1b5
--- /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});
+  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 b909d9c9392f..c7d37b943b65 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,22 @@
 
 #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>
+#include <stdexcept>
 
 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 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>>;
 using column      = cudf::column;
@@ -163,4 +170,130 @@ 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(
+    std::ignore = cudf::tables_equal(cudf::table_view{{left}}, cudf::table_view{{right}}),
+    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()
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);
 }