From 53bd4916d108f1a9a8e49de57856b9f54dfd7477 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Mon, 19 Oct 2026 17:19:41 -0700
Subject: [PATCH 1/2] Avoid reduction overhead in hash join probe kernels

Select separate sum and no-sum kernel specializations on the host so
match contexts and join retrieval do not reserve reduction resources.
Retain the fused reduction for size queries and shared probing logic.
---
 cpp/src/join/hash_join/hash_csr_kernels.cuh | 39 +++++++++++++--------
 1 file changed, 24 insertions(+), 15 deletions(-)

diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh
index ae577f547c32..c40f5ab9d47a 100644
--- a/cpp/src/join/hash_join/hash_csr_kernels.cuh
+++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh
@@ -127,7 +127,7 @@ CUDF_KERNEL void hash_csr_build_fill_kernel(size_type num_rows,
   }
 }
 
-template <bool IsOuter, typename Equal, typename Hasher>
+template <bool IsOuter, bool ComputeSum, typename Equal, typename Hasher>
 CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows,
                                              bitmask_type const* valid_rows,
                                              size_type* probe_groups,
@@ -156,7 +156,7 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows,
     }
     auto const match_count = IsOuter ? cuda::std::max(count, size_type{1}) : count;
     if (match_counts != nullptr) { match_counts[index] = match_count; }
-    thread_matches += static_cast<cuda::std::uint64_t>(match_count);
+    if constexpr (ComputeSum) { thread_matches += static_cast<cuda::std::uint64_t>(match_count); }
 
     // Only right and full joins consume the matched-row tally, and `matched_groups` is null for
     // every other kind, so this whole block compiles away outside outer joins rather than costing
@@ -175,7 +175,7 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows,
     }
   }
 
-  if (total_matches != nullptr) {
+  if constexpr (ComputeSum) {
     using block_reduce = cub::BlockReduce<cuda::std::uint64_t, hash_csr_block_size>;
     __shared__ typename block_reduce::TempStorage temp_storage;
     auto const block_matches = block_reduce(temp_storage).Sum(thread_matches);
@@ -251,18 +251,27 @@ void launch_hash_csr_probe_count_kernel(size_type num_rows,
 {
   if (num_rows == 0) { return; }
   auto const config = grid_1d{num_rows, hash_csr_block_size};
-  hash_csr_probe_count_kernel<IsOuter>
-    <<<config.num_blocks, config.num_threads_per_block, 0, stream.get()>>>(num_rows,
-                                                                           valid_rows,
-                                                                           probe_groups,
-                                                                           match_counts,
-                                                                           total_matches,
-                                                                           matched_groups,
-                                                                           matched_build_rows,
-                                                                           map,
-                                                                           csr,
-                                                                           equal,
-                                                                           hasher);
+  // Keep reduction registers and shared memory out of the kernels used by match contexts and
+  // retrieval. A device-side null check still reserves those resources for every probe.
+  auto const launch_kernel = [&]<bool ComputeSum>() {
+    hash_csr_probe_count_kernel<IsOuter, ComputeSum>
+      <<<config.num_blocks, config.num_threads_per_block, 0, stream.get()>>>(num_rows,
+                                                                             valid_rows,
+                                                                             probe_groups,
+                                                                             match_counts,
+                                                                             total_matches,
+                                                                             matched_groups,
+                                                                             matched_build_rows,
+                                                                             map,
+                                                                             csr,
+                                                                             equal,
+                                                                             hasher);
+  };
+  if (total_matches != nullptr) {
+    launch_kernel.template operator()<true>();
+  } else {
+    launch_kernel.template operator()<false>();
+  }
   CUDF_CUDA_TRY(cudaGetLastError());
 }
 

From 8c2fe8584a522c6d51ff6e5bb6943cd0fad3cb84 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Tue, 20 Oct 2026 10:29:52 -0700
Subject: [PATCH 2/2] Use bool dispatch helper for hash join probe launches

---
 cpp/src/join/hash_join/hash_csr_kernels.cuh | 12 ++++--------
 1 file changed, 4 insertions(+), 8 deletions(-)

diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh
index c40f5ab9d47a..1052605a3000 100644
--- a/cpp/src/join/hash_join/hash_csr_kernels.cuh
+++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh
@@ -10,6 +10,7 @@
 
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/cuda.hpp>
+#include <cudf/detail/utilities/dispatchers.hpp>
 #include <cudf/detail/utilities/grid_1d.cuh>
 #include <cudf/detail/utilities/integer_utils.hpp>
 #include <cudf/utilities/bit.hpp>
@@ -253,8 +254,8 @@ void launch_hash_csr_probe_count_kernel(size_type num_rows,
   auto const config = grid_1d{num_rows, hash_csr_block_size};
   // Keep reduction registers and shared memory out of the kernels used by match contexts and
   // retrieval. A device-side null check still reserves those resources for every probe.
-  auto const launch_kernel = [&]<bool ComputeSum>() {
-    hash_csr_probe_count_kernel<IsOuter, ComputeSum>
+  dispatch_bool(total_matches != nullptr, [&](auto compute_sum) {
+    hash_csr_probe_count_kernel<IsOuter, decltype(compute_sum)::value>
       <<<config.num_blocks, config.num_threads_per_block, 0, stream.get()>>>(num_rows,
                                                                              valid_rows,
                                                                              probe_groups,
@@ -266,12 +267,7 @@ void launch_hash_csr_probe_count_kernel(size_type num_rows,
                                                                              csr,
                                                                              equal,
                                                                              hasher);
-  };
-  if (total_matches != nullptr) {
-    launch_kernel.template operator()<true>();
-  } else {
-    launch_kernel.template operator()<false>();
-  }
+  });
   CUDF_CUDA_TRY(cudaGetLastError());
 }