Replace hash join internals with HashCSR

cpp/CMakeLists.txt

@@ -865,7 +865,6 @@ add_library(
   src/join/hash_join/full_join_match_context.cpp
   src/join/hash_join/full_join_retrieve.cu
   src/join/hash_join/full_join_size.cu
-  src/join/hash_join/full_join_size_impl.cu
   src/join/hash_join/hash_join.cu
   src/join/hash_join/inner_join_match_context.cpp
   src/join/hash_join/inner_join_retrieve.cu
@@ -874,14 +873,10 @@ add_library(
   src/join/hash_join/left_join_retrieve.cu
   src/join/hash_join/left_join_size.cu
   src/join/hash_join/match_context.cu
-  src/join/hash_join/partitioned_count.cu
-  src/join/hash_join/partitioned_count_outer.cu
   src/join/hash_join/partitioned_full_join.cu
   src/join/hash_join/partitioned_inner_join.cu
   src/join/hash_join/partitioned_join_retrieve.cu
   src/join/hash_join/partitioned_left_join.cu
-  src/join/hash_join/partitioned_retrieve.cu
-  src/join/hash_join/partitioned_retrieve_outer.cu
   src/join/join.cu
   src/join/join_utils.cu
   src/join/key_remapping.cu

cpp/include/cudf/detail/join/hash_join.hpp

@@ -181,7 +181,7 @@ class hash_join {
                         rmm::device_async_resource_ref mr) const;
 
  private:
-  bool const _is_empty;   ///< true if `_hash_table` is empty
+  bool const _is_empty;   ///< true if the build-side (right) table is empty
   bool const _has_nulls;  ///< true if nulls are present in either right table or any left table
   cudf::null_equality const _nulls_equal;  ///< whether to consider nulls as equal
   cudf::table_view _right;                 ///< input table to build the hash map

cpp/src/join/hash_join/common.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -20,22 +20,9 @@
 namespace cudf::detail {
 
 using hash_join_hasher = cudf::hashing::detail::MurmurHash3_x86_32<cudf::hash_value_type>;
-using hash_table_t     = typename cudf::detail::hash_join<hash_join_hasher>::impl::hash_table_t;
 
 bool is_trivial_join(table_view const& left, table_view const& right, join_kind join_type);
 
 void validate_hash_join_probe(table_view const& right, table_view const& left, bool has_nulls);
 
-std::unique_ptr<rmm::device_uvector<size_type>> make_join_match_counts(
-  table_view const& right,
-  std::shared_ptr<cudf::detail::row::equality::preprocessed_table> const& preprocessed_right,
-  cudf::detail::hash_table_t const& hash_table,
-  bool is_empty,
-  bool has_nulls,
-  null_equality compare_nulls,
-  join_kind join,
-  table_view const& left,
-  rmm::cuda_stream_view stream,
-  rmm::device_async_resource_ref mr);
-
 }  // namespace cudf::detail

cpp/src/join/hash_join/dispatch.cuh

@@ -1,12 +1,11 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
 
 #include "common.cuh"
 
-#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/row_operator/equality.cuh>
 #include <cudf/detail/row_operator/hashing.cuh>
 #include <cudf/detail/row_operator/primitive_row_operators.cuh>
@@ -41,17 +40,6 @@ class pair_equal {
   Equal _check_row_equality;
 };
 
-/**
- * @brief Extracts the right-side row index from a cuco hash table slot.
- */
-struct output_fn {
-  __device__ constexpr cudf::size_type operator()(
-    cuco::pair<hash_value_type, cudf::size_type> const& slot) const
-  {
-    return slot.second;
-  }
-};
-
 /**
  * @brief Equality comparator for cuco hash table probing with primitive row equality.
  */

cpp/src/join/hash_join/hash_csr.cuh

@@ -0,0 +1,66 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <cudf/detail/cuco_helpers.hpp>
+#include <cudf/hashing.hpp>
+#include <cudf/types.hpp>
+
+#include <cuco/pair.cuh>
+#include <cuda/atomic>
+
+#include <cstdint>
+
+namespace cudf::detail {
+
+using hash_csr_key_type            = cuco::pair<hash_value_type, size_type>;
+using hash_csr_build_position_type = cuco::pair<std::uint32_t, size_type>;
+
+struct hash_csr_map_view {
+  hash_csr_key_type* entries;
+  std::uint32_t capacity;
+  std::uint32_t mask;
+
+  template <typename Equal>
+  __device__ std::uint32_t insert(hash_csr_key_type key, Equal equal) const
+  {
+    for (std::uint32_t step = 0; step < capacity; ++step) {
+      auto const slot = (static_cast<std::uint32_t>(key.first) + step) & mask;
+      auto entry_ref =
+        cuda::atomic_ref<hash_csr_key_type, cuda::thread_scope_device>{entries[slot]};
+      auto old = hash_csr_key_type{hash_value_type{-1}, CUDF_SIZE_TYPE_SENTINEL};
+      if (entry_ref.compare_exchange_strong(old, key, cuda::memory_order_relaxed)) { return slot; }
+      if (equal(key, old)) { return slot; }
+    }
+    return capacity;
+  }
+
+  template <typename Equal>
+  __device__ std::uint32_t find(hash_csr_key_type key, Equal equal) const
+  {
+    for (std::uint32_t step = 0; step < capacity; ++step) {
+      auto const slot    = (static_cast<std::uint32_t>(key.first) + step) & mask;
+      auto const current = entries[slot];
+      if (current.second == CUDF_SIZE_TYPE_SENTINEL) { return capacity; }
+      if (equal(key, current)) { return slot; }
+    }
+    return capacity;
+  }
+};
+
+struct hash_csr_view {
+  size_type const* cumulative_ends;
+  size_type const* values;
+
+  __device__ size_type begin(size_type slot) const
+  {
+    return slot == 0 ? size_type{0} : cumulative_ends[slot - 1];
+  }
+
+  __device__ size_type size(size_type slot) const { return cumulative_ends[slot] - begin(slot); }
+};
+
+}  // namespace cudf::detail