From 6552766c91fbf15b2ab2f969f7d21b6250f743be Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 11 Aug 2026 19:26:55 +0000 Subject: [PATCH 01/12] Replace hash join internals with HashCSR --- cpp/include/cudf/detail/join/hash_join.hpp | 2 +- cpp/src/join/hash_join/common.cuh | 15 +- cpp/src/join/hash_join/dispatch.cuh | 73 +++- cpp/src/join/hash_join/full_join_size_impl.cu | 2 + cpp/src/join/hash_join/hash_csr.cuh | 104 +++++ cpp/src/join/hash_join/hash_csr_kernels.cuh | 360 ++++++++++++++++++ cpp/src/join/hash_join/hash_join.cu | 103 +++-- cpp/src/join/hash_join/hash_join_impl.cuh | 72 ++-- cpp/src/join/hash_join/match_context.cu | 93 ++--- cpp/src/join/hash_join/partitioned_count.cu | 4 +- .../join/hash_join/partitioned_count_outer.cu | 4 +- .../hash_join/partitioned_join_retrieve.cu | 129 ++++--- .../join/hash_join/partitioned_retrieve.cu | 4 +- .../hash_join/partitioned_retrieve_outer.cu | 4 +- cpp/src/join/hash_join/retrieve_impl.cuh | 219 ++++------- cpp/src/join/hash_join/size_impl.cuh | 136 ++++--- cpp/src/join/join.cu | 118 +++--- 17 files changed, 969 insertions(+), 473 deletions(-) create mode 100644 cpp/src/join/hash_join/hash_csr.cuh create mode 100644 cpp/src/join/hash_join/hash_csr_kernels.cuh diff --git a/cpp/include/cudf/detail/join/hash_join.hpp b/cpp/include/cudf/detail/join/hash_join.hpp index b3f5f8f5a39a..4a04234b8d70 100644 --- a/cpp/include/cudf/detail/join/hash_join.hpp +++ b/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 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 diff --git a/cpp/src/join/hash_join/common.cuh b/cpp/src/join/hash_join/common.cuh index d80f981099ef..1a6fec8c8ba2 100644 --- a/cpp/src/join/hash_join/common.cuh +++ b/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; -using hash_table_t = typename cudf::detail::hash_join::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> make_join_match_counts( - table_view const& right, - std::shared_ptr 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 diff --git a/cpp/src/join/hash_join/dispatch.cuh b/cpp/src/join/hash_join/dispatch.cuh index 2092817e2931..9a5006d6d286 100644 --- a/cpp/src/join/hash_join/dispatch.cuh +++ b/cpp/src/join/hash_join/dispatch.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 @@ -73,6 +73,77 @@ class primitive_pair_equal { cudf::detail::row::primitive::row_equality_comparator _check_row_equality; }; +template +class hash_csr_pair_equal { + public: + explicit hash_csr_pair_equal(Equal check_row_equality) + : _check_row_equality{std::move(check_row_equality)} + { + } + + __device__ __forceinline__ bool operator()(probe_key_type const& lhs, + probe_key_type const& rhs) const noexcept + { + using detail::row::lhs_index_type; + using detail::row::rhs_index_type; + return _check_row_equality(lhs_index_type{lhs.second}, rhs_index_type{rhs.second}); + } + + private: + Equal _check_row_equality; +}; + +class primitive_hash_csr_pair_equal { + public: + explicit primitive_hash_csr_pair_equal( + cudf::detail::row::primitive::row_equality_comparator check_row_equality) + : _check_row_equality{std::move(check_row_equality)} + { + } + + __device__ __forceinline__ bool operator()(probe_key_type const& lhs, + probe_key_type const& rhs) const noexcept + { + return _check_row_equality(lhs.second, rhs.second); + } + + private: + cudf::detail::row::primitive::row_equality_comparator _check_row_equality; +}; + +template +decltype(auto) dispatch_hash_csr_comparator( + table_view const& right_table, + table_view const& left_table, + std::shared_ptr const& preprocessed_right, + std::shared_ptr const& preprocessed_left, + bool has_nulls, + null_equality compare_nulls, + Fn&& fn) +{ + auto const left_nulls = cudf::nullate::DYNAMIC{has_nulls}; + + if (cudf::detail::is_primitive_row_op_compatible(right_table)) { + auto const d_hasher = cudf::detail::row::primitive::row_hasher{left_nulls, preprocessed_left}; + auto const d_equal = cudf::detail::row::primitive::row_equality_comparator{ + left_nulls, preprocessed_left, preprocessed_right, compare_nulls}; + return std::forward(fn)(primitive_hash_csr_pair_equal{d_equal}, d_hasher); + } + + auto const d_hasher = + cudf::detail::row::hash::row_hasher{preprocessed_left}.device_hasher(left_nulls); + auto const row_comparator = + cudf::detail::row::equality::two_table_comparator{preprocessed_left, preprocessed_right}; + + if (cudf::detail::has_nested_columns(left_table)) { + auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); + return std::forward(fn)(hash_csr_pair_equal{d_equal}, d_hasher); + } + + auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); + return std::forward(fn)(hash_csr_pair_equal{d_equal}, d_hasher); +} + template decltype(auto) dispatch_join_comparator( table_view const& right_table, diff --git a/cpp/src/join/hash_join/full_join_size_impl.cu b/cpp/src/join/hash_join/full_join_size_impl.cu index 0e9afcde668d..13bb1393f64e 100644 --- a/cpp/src/join/hash_join/full_join_size_impl.cu +++ b/cpp/src/join/hash_join/full_join_size_impl.cu @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#if 0 // Replaced by the HashCSR full-join size path in size_impl.cuh. #include "retrieve_impl.cuh" #include @@ -96,3 +97,4 @@ std::size_t get_full_join_size( } } // namespace cudf::detail +#endif diff --git a/cpp/src/join/hash_join/hash_csr.cuh b/cpp/src/join/hash_join/hash_csr.cuh new file mode 100644 index 000000000000..4a68b754f6cd --- /dev/null +++ b/cpp/src/join/hash_join/hash_csr.cuh @@ -0,0 +1,104 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include "kernels_common.cuh" + +#include + +#include +#include + +namespace cudf::detail { + +constexpr size_type hash_csr_empty_slot = size_type{-1}; +constexpr std::uint64_t hash_csr_empty_entry = std::numeric_limits::max(); +constexpr std::uint64_t hash_csr_empty_build_position = std::numeric_limits::max(); + +__device__ inline std::uint64_t pack_hash_csr_entry(hash_value_type hash, size_type row) +{ + return (static_cast(hash) << 32) | static_cast(row); +} + +__device__ inline hash_value_type unpack_hash_csr_hash(std::uint64_t value) +{ + return static_cast(value >> 32); +} + +__device__ inline size_type unpack_hash_csr_row(std::uint64_t value) +{ + return static_cast(static_cast(value)); +} + +__device__ inline std::uint64_t pack_hash_csr_build_position(std::uint32_t slot, size_type rank) +{ + return (static_cast(slot) << 32) | static_cast(rank); +} + +__device__ inline std::uint32_t unpack_hash_csr_build_slot(std::uint64_t value) +{ + return static_cast(value >> 32); +} + +__device__ inline size_type unpack_hash_csr_build_rank(std::uint64_t value) +{ + return static_cast(static_cast(value)); +} + +struct hash_csr_map_view { + std::uint64_t* entries; + std::uint32_t capacity; + std::uint32_t mask; + + template + __device__ std::uint32_t find_or_insert(probe_key_type key, Equal equal) const + { + auto const desired = pack_hash_csr_entry(key.first, key.second); + for (std::uint32_t step = 0; step < capacity; ++step) { + auto const slot = (static_cast(key.first) + step) & mask; + auto const old = atomicCAS(reinterpret_cast(entries + slot), + static_cast(hash_csr_empty_entry), + static_cast(desired)); + if (old == hash_csr_empty_entry) { return slot; } + if (unpack_hash_csr_hash(old) == key.first && + equal(key, probe_key_type{unpack_hash_csr_hash(old), unpack_hash_csr_row(old)})) { + return slot; + } + } + return capacity; + } + + template + __device__ std::uint32_t find(probe_key_type key, Equal equal) const + { + for (std::uint32_t step = 0; step < capacity; ++step) { + auto const slot = (static_cast(key.first) + step) & mask; + auto const current = entries[slot]; + if (current == hash_csr_empty_entry) { return capacity; } + if (unpack_hash_csr_hash(current) == key.first && + equal(key, probe_key_type{unpack_hash_csr_hash(current), unpack_hash_csr_row(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 end(size_type slot) const { return cumulative_ends[slot]; } + + __device__ size_type size(size_type slot) const { return end(slot) - begin(slot); } +}; + +} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh new file mode 100644 index 000000000000..67b7c4bad877 --- /dev/null +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -0,0 +1,360 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include "hash_csr.cuh" + +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include + +#include +#include +#include + +namespace cudf::detail { + +constexpr thread_index_type hash_csr_block_size = 256; +constexpr thread_index_type hash_csr_warp_size = 32; +constexpr thread_index_type hash_csr_warps_per_block = hash_csr_block_size / hash_csr_warp_size; +constexpr thread_index_type hash_csr_outputs_per_lane = 32; + +template +CUDF_KERNEL void hash_csr_build_count_kernel(size_type num_rows, + bitmask_type const* valid_rows, + std::uint64_t* build_positions, + size_type* slot_counts, + hash_csr_map_view map, + Equal equal, + Hasher hasher) +{ + auto const stride = grid_1d::grid_stride(); + for (auto row = grid_1d::global_thread_id(); row < num_rows; row += stride) { + auto const index = static_cast(row); + if (valid_rows != nullptr && !cudf::bit_is_set(valid_rows, index)) { + build_positions[index] = hash_csr_empty_build_position; + continue; + } + + auto const slot = map.find_or_insert(probe_key_type{hasher(index), index}, equal); + if (slot == map.capacity) { + build_positions[index] = hash_csr_empty_build_position; + continue; + } + auto const rank = atomicAdd(slot_counts + slot, size_type{1}); + build_positions[index] = pack_hash_csr_build_position(slot, rank); + } +} + +CUDF_KERNEL void hash_csr_build_fill_kernel(size_type num_rows, + std::uint64_t const* build_positions, + size_type const* cumulative_ends, + size_type* values) +{ + auto const stride = grid_1d::grid_stride(); + for (auto row = grid_1d::global_thread_id(); row < num_rows; row += stride) { + auto const index = static_cast(row); + auto const position = build_positions[index]; + if (position == hash_csr_empty_build_position) { continue; } + auto const slot = unpack_hash_csr_build_slot(position); + auto const rank = unpack_hash_csr_build_rank(position); + auto const begin = slot == 0 ? size_type{0} : cumulative_ends[slot - 1]; + values[begin + rank] = index; + } +} + +template +CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, + bitmask_type const* valid_rows, + size_type* probe_slots, + size_type* match_counts, + std::uint32_t* matched_slots, + unsigned long long* matched_build_rows, + hash_csr_map_view map, + hash_csr_view csr, + Equal equal, + Hasher hasher) +{ + auto const stride = grid_1d::grid_stride(); + for (auto row = grid_1d::global_thread_id(); row < num_rows; row += stride) { + auto const index = static_cast(row); + auto slot = map.capacity; + if (valid_rows == nullptr || cudf::bit_is_set(valid_rows, index)) { + slot = map.find(probe_key_type{hasher(index), index}, equal); + } + + auto const found = slot != map.capacity; + auto const count = found ? csr.size(static_cast(slot)) : size_type{0}; + if (probe_slots != nullptr) { + probe_slots[index] = found ? static_cast(slot) : hash_csr_empty_slot; + } + if (match_counts != nullptr) { + match_counts[index] = is_outer ? cuda::std::max(count, size_type{1}) : count; + } + + if (found && matched_slots != nullptr && + atomicCAS(matched_slots + slot, std::uint32_t{0}, std::uint32_t{1}) == 0) { + atomicAdd(matched_build_rows, static_cast(count)); + } + } +} + +template +void launch_hash_csr_build_count(size_type num_rows, + bitmask_type const* valid_rows, + std::uint64_t* build_positions, + size_type* slot_counts, + hash_csr_map_view map, + Equal equal, + Hasher hasher, + rmm::cuda_stream_view stream) +{ + if (num_rows == 0) { return; } + auto const config = grid_1d{num_rows, hash_csr_block_size}; + hash_csr_build_count_kernel<<>>( + num_rows, valid_rows, build_positions, slot_counts, map, equal, hasher); + CUDF_CUDA_TRY(cudaGetLastError()); +} + +inline void launch_hash_csr_build_fill(size_type num_rows, + std::uint64_t const* build_positions, + size_type const* cumulative_ends, + size_type* values, + rmm::cuda_stream_view stream) +{ + if (num_rows == 0) { return; } + auto const config = grid_1d{num_rows, hash_csr_block_size}; + hash_csr_build_fill_kernel<<>>( + num_rows, build_positions, cumulative_ends, values); + CUDF_CUDA_TRY(cudaGetLastError()); +} + +template +void launch_hash_csr_probe_count(size_type num_rows, + bitmask_type const* valid_rows, + size_type* probe_slots, + size_type* match_counts, + std::uint32_t* matched_slots, + unsigned long long* matched_build_rows, + hash_csr_map_view map, + hash_csr_view csr, + Equal equal, + Hasher hasher, + rmm::cuda_stream_view stream) +{ + if (num_rows == 0) { return; } + auto const config = grid_1d{num_rows, hash_csr_block_size}; + hash_csr_probe_count_kernel + <<>>(num_rows, + valid_rows, + probe_slots, + match_counts, + matched_slots, + matched_build_rows, + map, + csr, + equal, + hasher); + CUDF_CUDA_TRY(cudaGetLastError()); +} + +template +void hash_csr_exclusive_scan(InputIterator input, + OutputIterator output, + size_type num_items, + InitialValue initial_value, + rmm::cuda_stream_view stream) +{ + auto const mr = cudf::get_current_device_resource_ref(); + std::size_t temp_storage_bytes{}; + CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan(nullptr, + temp_storage_bytes, + input, + output, + cuda::std::plus<>{}, + initial_value, + num_items, + stream.value())); + rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); + CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan(temp_storage.data(), + temp_storage_bytes, + input, + output, + cuda::std::plus<>{}, + initial_value, + num_items, + stream.value())); +} + +inline void hash_csr_inclusive_sum(size_type* values, size_type size, rmm::cuda_stream_view stream) +{ + auto const mr = cudf::get_current_device_resource_ref(); + std::size_t temp_storage_bytes{}; + CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum( + nullptr, temp_storage_bytes, values, values, size, stream.value())); + rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); + CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum( + temp_storage.data(), temp_storage_bytes, values, values, size, stream.value())); +} + +inline std::int64_t hash_csr_scan_counts(size_type const* counts, + size_type num_rows, + std::int64_t* offsets, + rmm::cuda_stream_view stream) +{ + auto const mr = cudf::get_current_device_resource_ref(); + cudf::detail::device_scalar output_size(stream, mr); + auto output = cudf::detail::make_sizes_to_offsets_iterator( + offsets, offsets + num_rows + 1, output_size.data()); + hash_csr_exclusive_scan(counts, output, num_rows + 1, std::int64_t{0}, stream); + return output_size.value(stream); +} + +struct hash_csr_count_to_int64 { + __device__ std::int64_t operator()(size_type value) const + { + return static_cast(value); + } +}; + +inline std::int64_t hash_csr_reduce_counts(size_type const* counts, + size_type num_rows, + rmm::cuda_stream_view stream) +{ + if (num_rows == 0) { return 0; } + auto const mr = cudf::get_current_device_resource_ref(); + auto input = cuda::transform_iterator{counts, hash_csr_count_to_int64{}}; + cudf::detail::device_scalar result(stream, mr); + std::size_t temp_storage_bytes{}; + CUDF_CUDA_TRY(cub::DeviceReduce::Sum( + nullptr, temp_storage_bytes, input, result.data(), num_rows, stream.value())); + rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); + CUDF_CUDA_TRY(cub::DeviceReduce::Sum( + temp_storage.data(), temp_storage_bytes, input, result.data(), num_rows, stream.value())); + return result.value(stream); +} + +__device__ inline size_type hash_csr_find_probe_row_in_range(std::int64_t const* offsets, + size_type first_probe, + size_type last_probe, + std::int64_t output_index) +{ + auto first = static_cast(first_probe); + auto last = static_cast(last_probe) + 2; + while (first < last) { + auto const middle = first + (last - first) / 2; + if (offsets[middle] <= output_index) { + first = middle + 1; + } else { + last = middle; + } + } + return static_cast(first - 1); +} + +template +CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, + size_type num_probe_rows, + std::int64_t outputs_per_warp, + std::int64_t const* offsets, + size_type const* probe_slots, + hash_csr_view csr, + size_type left_index_offset, + size_type* left_indices, + size_type* right_indices) +{ + auto const lane_id = static_cast(threadIdx.x) % hash_csr_warp_size; + auto const warp_in_block = static_cast(threadIdx.x) / hash_csr_warp_size; + auto const global_warp = + static_cast(blockIdx.x) * hash_csr_warps_per_block + warp_in_block; + auto const range_begin = outputs_per_warp * global_warp; + if (range_begin >= output_size) { return; } + auto const range_end = cuda::std::min(range_begin + outputs_per_warp, output_size); + + size_type endpoint_probe{}; + if (lane_id < 2) { + auto const endpoint = lane_id == 0 ? range_begin : range_end - 1; + endpoint_probe = hash_csr_find_probe_row_in_range(offsets, 0, num_probe_rows - 1, endpoint); + } + auto const first_probe = __shfl_sync(0xffffffff, endpoint_probe, 0); + auto const last_probe = __shfl_sync(0xffffffff, endpoint_probe, 1); + +#pragma unroll + for (thread_index_type item = 0; item < hash_csr_outputs_per_lane; ++item) { + auto const output_index = range_begin + lane_id + item * hash_csr_warp_size; + if (output_index < range_end) { + auto const probe_row = + first_probe == last_probe + ? first_probe + : hash_csr_find_probe_row_in_range(offsets, first_probe, last_probe, output_index); + auto const slot = probe_slots[probe_row]; + left_indices[output_index] = probe_row + left_index_offset; + if constexpr (is_outer) { + if (slot == hash_csr_empty_slot) { + right_indices[output_index] = JoinNoMatch; + continue; + } + } + auto const local_match = static_cast(output_index - offsets[probe_row]); + right_indices[output_index] = csr.values[csr.begin(slot) + local_match]; + } + } +} + +template +void launch_hash_csr_retrieve(std::int64_t output_size, + size_type num_probe_rows, + std::int64_t const* offsets, + size_type const* probe_slots, + hash_csr_view csr, + size_type left_index_offset, + size_type* left_indices, + size_type* right_indices, + rmm::cuda_stream_view stream) +{ + if (output_size == 0) { return; } + auto const min_blocks = size_type{2} * cudf::detail::num_multiprocessors(); + constexpr auto outputs_per_block = + hash_csr_warps_per_block * hash_csr_warp_size * hash_csr_outputs_per_lane; + auto const requested_blocks = (output_size + outputs_per_block - 1) / outputs_per_block; + auto const num_blocks = + static_cast(cuda::std::max(requested_blocks, min_blocks)); + auto const num_warps = static_cast(num_blocks) * hash_csr_warps_per_block; + auto const outputs_per_warp = (output_size + num_warps - 1) / num_warps; + + hash_csr_retrieve_kernel + <<>>(output_size, + num_probe_rows, + outputs_per_warp, + offsets, + probe_slots, + csr, + left_index_offset, + left_indices, + right_indices); + CUDF_CUDA_TRY(cudaGetLastError()); +} + +} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index b699e04fff42..027381ee817d 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -4,6 +4,8 @@ */ #include "common.cuh" +#include "dispatch.cuh" +#include "hash_csr_kernels.cuh" #include "join/join_common_utils.cuh" #include @@ -23,8 +25,11 @@ #include +#include +#include #include #include +#include #include namespace cudf::detail { @@ -59,43 +64,19 @@ void validate_hash_join_probe(table_view const& right, table_view const& left, b } namespace { -void build_hash_join( - cudf::table_view const& right, - std::shared_ptr const& preprocessed_right, - cudf::detail::hash_table_t& hash_table, - bool has_nested_nulls, - null_equality nulls_equal, - [[maybe_unused]] bitmask_type const* bitmask, - rmm::cuda_stream_view stream) +std::uint32_t hash_csr_capacity(size_type rows, double load_factor) { - CUDF_EXPECTS(0 != right.num_columns(), "Selected right dataset is empty", std::invalid_argument); - CUDF_EXPECTS(0 != right.num_rows(), "Right side table has no rows", std::invalid_argument); - - auto insert_rows = [&](auto const& right, auto const& d_hasher) { - auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_fn{d_hasher}); - - if (nulls_equal == cudf::null_equality::EQUAL or not nullable(right)) { - hash_table.insert(iter, iter + right.num_rows(), stream.value()); - } else { - auto const stencil = cuda::counting_iterator{0}; - auto const pred = row_is_valid{bitmask}; - - hash_table.insert_if(iter, iter + right.num_rows(), stencil, pred, stream.value()); - } - }; - - auto const nulls = nullate::DYNAMIC{has_nested_nulls}; - - if (cudf::detail::is_primitive_row_op_compatible(right)) { - auto const d_hasher = cudf::detail::row::primitive::row_hasher{nulls, preprocessed_right}; - - insert_rows(right, d_hasher); - } else { - auto const row_hash = detail::row::hash::row_hasher{preprocessed_right}; - auto const d_hasher = row_hash.device_hasher(nulls); - - insert_rows(right, d_hasher); + auto const checked = checked_load_factor(load_factor); + auto const requested = std::max(static_cast(rows) + 1, + std::ceil(static_cast(rows) / checked)); + std::uint64_t capacity = 1; + while (static_cast(capacity) < requested) { + capacity <<= 1; } + CUDF_EXPECTS(capacity <= std::numeric_limits::max(), + "HashCSR table capacity is not representable", + std::overflow_error); + return static_cast(capacity); } } // namespace @@ -119,32 +100,46 @@ hash_join::hash_join(cudf::table_view const& right, : _has_nulls(has_nulls), _is_empty{right.num_rows() == 0}, _nulls_equal{compare_nulls}, - _impl{std::make_unique(impl{typename impl::hash_table_t{ - cuco::extent{static_cast(right.num_rows())}, - checked_load_factor(load_factor), - cuco::empty_key{cuco::pair{std::numeric_limits::max(), cudf::JoinNoMatch}}, - {}, - {}, - {}, - {}, - rmm::mr::polymorphic_allocator{std::move(mr)}, - stream.value()}})}, _right{right}, - _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create(_right, stream)} + _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create(_right, stream)}, + _impl{std::make_unique( + hash_csr_capacity(right.num_rows(), load_factor), right.num_rows(), stream, mr)} { CUDF_FUNC_RANGE(); CUDF_EXPECTS(0 != right.num_columns(), "Hash join right table is empty", std::invalid_argument); if (_is_empty) { return; } - auto const row_bitmask = - cudf::detail::bitmask_and(right, stream, cudf::get_current_device_resource_ref()).first; - cudf::detail::build_hash_join(_right, - _preprocessed_right, - _impl->_hash_table, - _has_nulls, - _nulls_equal, - reinterpret_cast(row_bitmask.data()), + CUDF_CUDA_TRY(cudaMemsetAsync( + _impl->entries.data(), 0xff, _impl->entries.size() * sizeof(std::uint64_t), stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(_impl->cumulative_ends.data(), + 0, + _impl->cumulative_ends.size() * sizeof(size_type), + stream.value())); + + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const row_bitmask = cudf::detail::bitmask_and(right, stream, temp_mr).first; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL + ? static_cast(row_bitmask.data()) + : nullptr; + rmm::device_uvector build_positions(right.num_rows(), stream, temp_mr); + auto build = [&](auto equality, auto hasher) { + launch_hash_csr_build_count(right.num_rows(), + valid_rows, + build_positions.data(), + _impl->cumulative_ends.data(), + _impl->map_view(), + equality, + hasher, stream); + }; + dispatch_hash_csr_comparator( + right, right, _preprocessed_right, _preprocessed_right, _has_nulls, _nulls_equal, build); + hash_csr_inclusive_sum(_impl->cumulative_ends.data(), _impl->capacity, stream); + launch_hash_csr_build_fill(right.num_rows(), + build_positions.data(), + _impl->cumulative_ends.data(), + _impl->values.data(), + stream); } template hash_join::hash_join( diff --git a/cpp/src/join/hash_join/hash_join_impl.cuh b/cpp/src/join/hash_join/hash_join_impl.cuh index 18a65be10f76..592b953d3211 100644 --- a/cpp/src/join/hash_join/hash_join_impl.cuh +++ b/cpp/src/join/hash_join/hash_join_impl.cuh @@ -1,62 +1,44 @@ /* - * 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 "hash_csr.cuh" + #include #include -#include +#include -#include -#include +#include namespace cudf::detail { template struct hash_join::impl { - struct always_not_equal { - __device__ constexpr bool operator()( - cuco::pair const&, - cuco::pair const&) const noexcept - { - // multiset always insert - return false; - } - }; - - struct hasher1 { - __device__ constexpr hash_value_type operator()( - cuco::pair const& key) const noexcept - { - return key.first; - } - }; - - struct hasher2 { - hasher2(hash_value_type seed) : _hash{seed} {} - - __device__ constexpr hash_value_type operator()( - cuco::pair const& key) const noexcept - { - return _hash(key.first); - } - - private: - Hasher _hash; - }; - - using hash_table_t = - cuco::static_multiset, - cuco::extent, - cuda::thread_scope_device, - always_not_equal, - cuco::double_hashing, - rmm::mr::polymorphic_allocator, - cuco::storage<2>>; - - hash_table_t _hash_table; + impl(std::uint32_t capacity, + size_type rows, + rmm::cuda_stream_view stream, + cuda::mr::any_resource const& mr) + : entries(capacity, stream, mr), + cumulative_ends(capacity, stream, mr), + values(rows, stream, mr), + capacity{capacity} + { + } + + hash_csr_map_view map_view() const + { + return {const_cast(entries.data()), capacity, capacity - 1}; + } + + hash_csr_view csr_view() const { return {cumulative_ends.data(), values.data()}; } + + rmm::device_uvector entries; + rmm::device_uvector cumulative_ends; + rmm::device_uvector values; + std::uint32_t capacity; }; } // namespace cudf::detail diff --git a/cpp/src/join/hash_join/match_context.cu b/cpp/src/join/hash_join/match_context.cu index 5fc2dd5ba9cf..11b42848a381 100644 --- a/cpp/src/join/hash_join/match_context.cu +++ b/cpp/src/join/hash_join/match_context.cu @@ -1,39 +1,32 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "common.cuh" #include "dispatch.cuh" -#include "join/join_common_utils.cuh" -#include "partitioned_count_kernels.hpp" +#include "hash_csr_kernels.cuh" -#include +#include +#include #include #include -#include #include -#include namespace cudf::detail { -std::unique_ptr> make_join_match_counts( - table_view const& right, - std::shared_ptr const& preprocessed_right, - cudf::detail::hash_table_t const& hash_table, - bool is_empty, - bool has_nulls, - null_equality compare_nulls, +template +std::unique_ptr> hash_join::make_match_counts( join_kind join, table_view const& left, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::device_async_resource_ref mr) const { auto match_counts = std::make_unique>(left.num_rows(), stream, mr); - if (is_empty) { + if (_is_empty) { thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), match_counts->begin(), match_counts->end(), @@ -41,60 +34,52 @@ std::unique_ptr> make_join_match_counts( return match_counts; } - CUDF_EXPECTS(has_nulls || !cudf::has_nested_nulls(left), + CUDF_EXPECTS(_has_nulls || !cudf::has_nested_nulls(left), "Left table has nulls while right table was not hashed with null check.", std::invalid_argument); auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); - auto const left_table_num_rows = left.num_rows(); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL + ? static_cast(row_bitmask.data()) + : nullptr; - auto count_matches = [&](auto equality, auto d_hasher) { - // Precompute left keys: {hash(row_idx), row_idx} for each left row. - auto const n = static_cast(left_table_num_rows); - rmm::device_uvector left_keys(n, stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator(0), - cuda::counting_iterator(left_table_num_rows), - left_keys.begin(), - pair_fn{d_hasher}); - - auto const ref = hash_table.ref(cuco::op::count) - .rebind_key_eq(equality) - .rebind_hash_function(hash_table.hash_function()); + auto count_matches = [&](auto equality, auto hasher) { if (join == join_kind::INNER_JOIN) { - launch_partitioned_count(left_keys.data(), n, match_counts->begin(), ref, stream); + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + nullptr, + match_counts->data(), + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); } else { - // IsOuter=true handles the clamp (zero → 1) for LEFT/FULL joins internally. - launch_partitioned_count(left_keys.data(), n, match_counts->begin(), ref, stream); + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + nullptr, + match_counts->data(), + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); } }; - dispatch_join_comparator( - right, left, preprocessed_right, preprocessed_left, has_nulls, compare_nulls, count_matches); + dispatch_hash_csr_comparator( + _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); return match_counts; } -template -std::unique_ptr> hash_join::make_match_counts( - join_kind join, - cudf::table_view const& left, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) const -{ - return make_join_match_counts(_right, - _preprocessed_right, - _impl->_hash_table, - _is_empty, - _has_nulls, - _nulls_equal, - join, - left, - stream, - mr); -} - template std::unique_ptr> hash_join::make_match_counts(join_kind, cudf::table_view const&, diff --git a/cpp/src/join/hash_join/partitioned_count.cu b/cpp/src/join/hash_join/partitioned_count.cu index 67af2bf0d056..345f9f9a7053 100644 --- a/cpp/src/join/hash_join/partitioned_count.cu +++ b/cpp/src/join/hash_join/partitioned_count.cu @@ -1,8 +1,9 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ +#if 0 // Replaced by the HashCSR match-context count path. #include "partitioned_count_kernels.cuh" #include "ref_types.cuh" @@ -21,3 +22,4 @@ template void launch_partitioned_count( probe_key_type const*, thread_index_type, size_type*, flat_count_ref_t, rmm::cuda_stream_view); } // namespace cudf::detail +#endif diff --git a/cpp/src/join/hash_join/partitioned_count_outer.cu b/cpp/src/join/hash_join/partitioned_count_outer.cu index 3f4204cb1fa8..d12c775e2221 100644 --- a/cpp/src/join/hash_join/partitioned_count_outer.cu +++ b/cpp/src/join/hash_join/partitioned_count_outer.cu @@ -1,8 +1,9 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ +#if 0 // Replaced by the shared HashCSR outer count kernel. #include "partitioned_count_kernels.cuh" #include "ref_types.cuh" @@ -21,3 +22,4 @@ template void launch_partitioned_count( probe_key_type const*, thread_index_type, size_type*, flat_count_ref_t, rmm::cuda_stream_view); } // namespace cudf::detail +#endif diff --git a/cpp/src/join/hash_join/partitioned_join_retrieve.cu b/cpp/src/join/hash_join/partitioned_join_retrieve.cu index 77dcae1e9325..9c0d3b365624 100644 --- a/cpp/src/join/hash_join/partitioned_join_retrieve.cu +++ b/cpp/src/join/hash_join/partitioned_join_retrieve.cu @@ -1,19 +1,22 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "common.cuh" #include "dispatch.cuh" -#include "join/join_common_utils.cuh" +#include "hash_csr_kernels.cuh" #include "join/join_common_utils.hpp" -#include "partitioned_retrieve_kernels.hpp" #include +#include #include +#include +#include #include #include #include +#include #include #include @@ -21,7 +24,6 @@ #include #include #include -#include namespace cudf::detail { namespace { @@ -109,49 +111,90 @@ hash_join::partitioned_join_retrieve(join_kind join, auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left_partition_view, stream); - // For FULL_JOIN, probe with LEFT_JOIN semantics (no complement here) - bool const is_outer = (join != join_kind::INNER_JOIN); - - // launch_partitioned_retrieve reduces match counts to compute output size - // (total = last_offset + last_count), allocates output buffers, and launches the kernel. - auto const* partition_counts = match_ctx._match_counts->data() + left_start_idx; - auto const n = static_cast(partition_size); - - std::pair>, - std::unique_ptr>> - join_indices; - - auto retrieve_partition = [&](auto equality, auto d_hasher) { - // Precompute left keys for this partition slice. - rmm::device_uvector left_keys(n, stream); - thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator(0), - cuda::counting_iterator(partition_size), - left_keys.begin(), - pair_fn{d_hasher}); - - auto const ref = _impl->_hash_table.ref(cuco::op::count) - .rebind_key_eq(equality) - .rebind_hash_function(_impl->_hash_table.hash_function()); - - if (is_outer) { - join_indices = launch_partitioned_retrieve( - left_keys.data(), n, partition_counts, ref, left_start_idx, stream, mr); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto counts = cudf::detail::make_zeroed_device_uvector_async( + static_cast(partition_size) + 1, stream, temp_mr); + CUDF_CUDA_TRY( + cudf::detail::memcpy_async(counts.data(), + match_ctx._match_counts->data() + left_start_idx, + static_cast(partition_size) * sizeof(size_type), + stream)); + auto offsets = cudf::detail::make_zeroed_device_uvector_async( + static_cast(partition_size) + 1, stream, temp_mr); + auto const output_size = + hash_csr_scan_counts(counts.data(), partition_size, offsets.data(), stream); + CUDF_EXPECTS(output_size >= 0, "Join output size overflowed", std::overflow_error); + + rmm::device_uvector probe_slots(partition_size, stream, temp_mr); + auto const row_bitmask = cudf::detail::bitmask_and(left_partition_view, stream, temp_mr).first; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL + ? static_cast(row_bitmask.data()) + : nullptr; + auto save_slots = [&](auto equality, auto hasher) { + if (join == join_kind::INNER_JOIN) { + launch_hash_csr_probe_count(partition_size, + valid_rows, + probe_slots.data(), + nullptr, + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); } else { - join_indices = launch_partitioned_retrieve( - left_keys.data(), n, partition_counts, ref, left_start_idx, stream, mr); + launch_hash_csr_probe_count(partition_size, + valid_rows, + probe_slots.data(), + nullptr, + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); } }; + dispatch_hash_csr_comparator(_right, + left_partition_view, + _preprocessed_right, + preprocessed_left, + _has_nulls, + _nulls_equal, + save_slots); + + auto left_indices = std::make_unique>( + static_cast(output_size), stream, mr); + auto right_indices = std::make_unique>( + static_cast(output_size), stream, mr); + cudf::prefetch::detail::prefetch(*left_indices, stream); + cudf::prefetch::detail::prefetch(*right_indices, stream); + + if (join == join_kind::INNER_JOIN) { + launch_hash_csr_retrieve(output_size, + partition_size, + offsets.data(), + probe_slots.data(), + _impl->csr_view(), + left_start_idx, + left_indices->data(), + right_indices->data(), + stream); + } else { + launch_hash_csr_retrieve(output_size, + partition_size, + offsets.data(), + probe_slots.data(), + _impl->csr_view(), + left_start_idx, + left_indices->data(), + right_indices->data(), + stream); + } - dispatch_join_comparator(_right, - left_partition_view, - _preprocessed_right, - preprocessed_left, - _has_nulls, - _nulls_equal, - retrieve_partition); - - return join_indices; + return {std::move(left_indices), std::move(right_indices)}; } template std::pair>, diff --git a/cpp/src/join/hash_join/partitioned_retrieve.cu b/cpp/src/join/hash_join/partitioned_retrieve.cu index efb200f4c899..29ba4bc9e0b0 100644 --- a/cpp/src/join/hash_join/partitioned_retrieve.cu +++ b/cpp/src/join/hash_join/partitioned_retrieve.cu @@ -1,8 +1,9 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ +#if 0 // Replaced by the HashCSR partitioned pull path. #include "partitioned_retrieve_kernels.cuh" #include "ref_types.cuh" @@ -39,3 +40,4 @@ launch_partitioned_retrieve(probe_key_type const*, rmm::device_async_resource_ref); } // namespace cudf::detail +#endif diff --git a/cpp/src/join/hash_join/partitioned_retrieve_outer.cu b/cpp/src/join/hash_join/partitioned_retrieve_outer.cu index 68b8f591ada9..1f6a519d78ba 100644 --- a/cpp/src/join/hash_join/partitioned_retrieve_outer.cu +++ b/cpp/src/join/hash_join/partitioned_retrieve_outer.cu @@ -1,8 +1,9 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ +#if 0 // Replaced by the shared HashCSR outer pull kernel. #include "partitioned_retrieve_kernels.cuh" #include "ref_types.cuh" @@ -39,3 +40,4 @@ launch_partitioned_retrieve(probe_key_type const*, rmm::device_async_resource_ref); } // namespace cudf::detail +#endif diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index 5efe69afe850..d98507546216 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -6,146 +6,21 @@ #include "common.cuh" #include "dispatch.cuh" -#include "join/join_common_utils.cuh" +#include "hash_csr_kernels.cuh" #include "join/join_common_utils.hpp" -#include "size_impl.cuh" -#include +#include #include +#include #include #include #include #include #include -#include - -#include -#include namespace cudf::detail { -template -std::pair>, - std::unique_ptr>> -probe_join_hash_table( - cudf::table_view const& right_table, - cudf::table_view const& left_table, - std::shared_ptr const& preprocessed_right, - std::shared_ptr const& preprocessed_left, - cudf::detail::hash_table_t const& hash_table, - bool has_nulls, - null_equality compare_nulls, - std::optional output_size, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - static_assert(Join == join_kind::INNER_JOIN || Join == join_kind::LEFT_JOIN || - Join == join_kind::FULL_JOIN); - - constexpr auto size_join = Join == join_kind::FULL_JOIN ? join_kind::LEFT_JOIN : Join; - - std::size_t const join_size = output_size - ? *output_size - : compute_join_output_size(right_table, - left_table, - preprocessed_right, - preprocessed_left, - hash_table, - has_nulls, - compare_nulls, - stream); - - if (join_size == 0) { - return std::pair(std::make_unique>(0, stream, mr), - std::make_unique>(0, stream, mr)); - } - - auto left_indices = std::make_unique>(join_size, stream, mr); - auto right_indices = std::make_unique>(join_size, stream, mr); - cudf::prefetch::detail::prefetch(*left_indices, stream); - cudf::prefetch::detail::prefetch(*right_indices, stream); - - auto const left_table_num_rows = left_table.num_rows(); - auto const out_probe_begin = - cuda::make_transform_output_iterator(left_indices->begin(), output_fn{}); - auto const out_build_begin = - cuda::make_transform_output_iterator(right_indices->begin(), output_fn{}); - - auto retrieve_results = [&](auto equality, auto d_hasher) { - auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_fn{d_hasher}); - if constexpr (Join == join_kind::INNER_JOIN) { - hash_table.retrieve(iter, - iter + left_table_num_rows, - equality, - hash_table.hash_function(), - out_probe_begin, - out_build_begin, - stream.value()); - } else { - [[maybe_unused]] auto out_probe_end = hash_table - .retrieve_outer(iter, - iter + left_table_num_rows, - equality, - hash_table.hash_function(), - out_probe_begin, - out_build_begin, - stream.value()) - .first; - - if constexpr (Join == join_kind::FULL_JOIN) { - auto const actual_size = cuda::std::distance(out_probe_begin, out_probe_end); - left_indices->resize(actual_size, stream); - right_indices->resize(actual_size, stream); - } - } - }; - - dispatch_join_comparator(right_table, - left_table, - preprocessed_right, - preprocessed_left, - has_nulls, - compare_nulls, - retrieve_results); - - return std::pair(std::move(left_indices), std::move(right_indices)); -} - -template -void retrieve_left_join_build_indices( - cudf::table_view const& right_table, - cudf::table_view const& left_table, - std::shared_ptr const& preprocessed_right, - std::shared_ptr const& preprocessed_left, - cudf::detail::hash_table_t const& hash_table, - bool has_nulls, - null_equality compare_nulls, - RightOutputIterator out_build_begin, - rmm::cuda_stream_view stream) -{ - auto const left_table_num_rows = left_table.num_rows(); - - auto retrieve_results = [&](auto equality, auto d_hasher) { - auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_fn{d_hasher}); - hash_table.retrieve_outer(iter, - iter + left_table_num_rows, - equality, - hash_table.hash_function(), - cuda::make_discard_iterator(), - out_build_begin, - stream.value()); - }; - - dispatch_join_comparator(right_table, - left_table, - preprocessed_right, - preprocessed_left, - has_nulls, - compare_nulls, - retrieve_results); -} - template template std::pair>, @@ -176,16 +51,84 @@ hash_join::join_retrieve(cudf::table_view const& left, auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); - auto join_indices = cudf::detail::probe_join_hash_table(_right, - left, - _preprocessed_right, - preprocessed_left, - _impl->_hash_table, - _has_nulls, - _nulls_equal, - output_size, - stream, - mr); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto match_counts = cudf::detail::make_zeroed_device_uvector_async( + static_cast(left.num_rows()) + 1, stream, temp_mr); + rmm::device_uvector probe_slots(left.num_rows(), stream, temp_mr); + auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL + ? static_cast(row_bitmask.data()) + : nullptr; + + auto count_matches = [&](auto equality, auto hasher) { + if constexpr (Join == join_kind::INNER_JOIN) { + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + probe_slots.data(), + match_counts.data(), + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); + } else { + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + probe_slots.data(), + match_counts.data(), + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); + } + }; + dispatch_hash_csr_comparator( + _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); + + auto offsets = cudf::detail::make_zeroed_device_uvector_async( + static_cast(left.num_rows()) + 1, stream, temp_mr); + auto const actual_size = + hash_csr_scan_counts(match_counts.data(), left.num_rows(), offsets.data(), stream); + CUDF_EXPECTS(actual_size >= 0, "Join output size overflowed", std::overflow_error); + auto const join_size = Join != join_kind::FULL_JOIN && output_size.has_value() + ? *output_size + : static_cast(actual_size); + CUDF_EXPECTS(join_size == static_cast(actual_size), + "The provided join output size is incorrect"); + + auto left_indices = std::make_unique>(join_size, stream, mr); + auto right_indices = std::make_unique>(join_size, stream, mr); + cudf::prefetch::detail::prefetch(*left_indices, stream); + cudf::prefetch::detail::prefetch(*right_indices, stream); + + if constexpr (Join == join_kind::INNER_JOIN) { + launch_hash_csr_retrieve(actual_size, + left.num_rows(), + offsets.data(), + probe_slots.data(), + _impl->csr_view(), + 0, + left_indices->data(), + right_indices->data(), + stream); + } else { + launch_hash_csr_retrieve(actual_size, + left.num_rows(), + offsets.data(), + probe_slots.data(), + _impl->csr_view(), + 0, + left_indices->data(), + right_indices->data(), + stream); + } + + auto join_indices = std::pair(std::move(left_indices), std::move(right_indices)); if constexpr (Join == join_kind::FULL_JOIN) { return detail::finalize_full_join( diff --git a/cpp/src/join/hash_join/size_impl.cuh b/cpp/src/join/hash_join/size_impl.cuh index 3e20ebc7367e..972ba8ba0a9f 100644 --- a/cpp/src/join/hash_join/size_impl.cuh +++ b/cpp/src/join/hash_join/size_impl.cuh @@ -1,66 +1,19 @@ /* - * 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 "dispatch.cuh" -#include "join/join_common_utils.cuh" +#include "hash_csr_kernels.cuh" +#include #include +#include namespace cudf::detail { -std::size_t get_full_join_size( - cudf::table_view const& right_table, - cudf::table_view const& left_table, - std::shared_ptr const& preprocessed_right, - std::shared_ptr const& preprocessed_left, - cudf::detail::hash_table_t const& hash_table, - bool has_nulls, - null_equality compare_nulls, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); - -template -std::size_t compute_join_output_size( - table_view const& right_table, - table_view const& left_table, - std::shared_ptr const& preprocessed_right, - std::shared_ptr const& preprocessed_left, - cudf::detail::hash_table_t const& hash_table, - bool has_nulls, - cudf::null_equality nulls_equal, - rmm::cuda_stream_view stream) -{ - static_assert(Join == join_kind::INNER_JOIN || Join == join_kind::LEFT_JOIN); - - if (right_table.num_rows() == 0) { - return Join == join_kind::INNER_JOIN ? 0 : left_table.num_rows(); - } - - auto const left_table_num_rows = left_table.num_rows(); - - return dispatch_join_comparator( - right_table, - left_table, - preprocessed_right, - preprocessed_left, - has_nulls, - nulls_equal, - [&](auto equality, auto d_hasher) { - auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_fn{d_hasher}); - if constexpr (Join == join_kind::LEFT_JOIN) { - return hash_table.count_outer( - iter, iter + left_table_num_rows, equality, hash_table.hash_function(), stream.value()); - } else { - return hash_table.count( - iter, iter + left_table_num_rows, equality, hash_table.hash_function(), stream.value()); - } - }); -} - template template std::size_t hash_join::join_size(cudf::table_view const& left, @@ -83,14 +36,32 @@ std::size_t hash_join::join_size(cudf::table_view const& left, auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); - return cudf::detail::compute_join_output_size(_right, - left, - _preprocessed_right, - preprocessed_left, - _impl->_hash_table, - _has_nulls, - _nulls_equal, - stream); + auto const temp_mr = cudf::get_current_device_resource_ref(); + auto match_counts = + cudf::detail::make_zeroed_device_uvector_async(left.num_rows(), stream, temp_mr); + auto const row_bitmask = cudf::detail::bitmask_and(left, stream, temp_mr).first; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL + ? static_cast(row_bitmask.data()) + : nullptr; + + auto count_matches = [&](auto equality, auto hasher) { + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + nullptr, + match_counts.data(), + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); + }; + dispatch_hash_csr_comparator( + _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); + auto const output_size = hash_csr_reduce_counts(match_counts.data(), left.num_rows(), stream); + CUDF_EXPECTS(output_size >= 0, "Join output size overflowed", std::overflow_error); + return static_cast(output_size); } template @@ -102,7 +73,6 @@ std::size_t hash_join::join_size(cudf::table_view const& left, static_assert(Join == join_kind::FULL_JOIN); CUDF_FUNC_RANGE(); - if (_is_empty) { return left.num_rows(); } CUDF_EXPECTS(_has_nulls || !cudf::has_nested_nulls(left), @@ -111,16 +81,42 @@ std::size_t hash_join::join_size(cudf::table_view const& left, auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); - - return cudf::detail::get_full_join_size(_right, - left, - _preprocessed_right, - preprocessed_left, - _impl->_hash_table, - _has_nulls, - _nulls_equal, - stream, - mr); + auto match_counts = + cudf::detail::make_zeroed_device_uvector_async(left.num_rows(), stream, mr); + auto matched_slots = + cudf::detail::make_zeroed_device_uvector_async(_impl->capacity, stream, mr); + auto matched_build_rows = cudf::detail::device_scalar(0, stream, mr); + auto const row_bitmask = cudf::detail::bitmask_and(left, stream, mr).first; + auto const valid_rows = _nulls_equal == null_equality::UNEQUAL + ? static_cast(row_bitmask.data()) + : nullptr; + + auto count_matches = [&](auto equality, auto hasher) { + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + nullptr, + match_counts.data(), + matched_slots.data(), + matched_build_rows.data(), + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); + }; + dispatch_hash_csr_comparator( + _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); + + auto const left_output_size = + hash_csr_reduce_counts(match_counts.data(), left.num_rows(), stream); + auto const matched_right_rows = matched_build_rows.value(stream); + CUDF_EXPECTS(left_output_size >= 0, "Join output size overflowed", std::overflow_error); + auto const output_size = static_cast(left_output_size) + + static_cast(_right.num_rows()) - matched_right_rows; + CUDF_EXPECTS(output_size <= std::numeric_limits::max(), + "Join output size overflowed", + std::overflow_error); + return static_cast(output_size); } } // namespace cudf::detail diff --git a/cpp/src/join/join.cu b/cpp/src/join/join.cu index 9334125e05df..f8f788d2b848 100644 --- a/cpp/src/join/join.cu +++ b/cpp/src/join/join.cu @@ -1,11 +1,12 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "join_common_utils.hpp" #include #include +#include #include #include #include @@ -19,11 +20,21 @@ #include #include +#include #include namespace cudf { namespace detail { +namespace { +bool has_dictionary_columns(table_view const& table) +{ + return std::any_of(table.begin(), table.end(), [](column_view const& column) { + return column.type().id() == type_id::DICTIONARY32; + }); +} +} // namespace + std::pair>, std::unique_ptr>> inner_join(table_view const& left_input, @@ -32,31 +43,34 @@ inner_join(table_view const& left_input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - // Make sure any dictionary columns have matched key sets. - // This will return any new dictionary columns created as well as updated table_views. - auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, - stream, - cudf::get_current_device_resource_ref()); // temporary objects returned - - // now rebuild the table views with the updated ones - auto const left = matched.second.front(); - auto const right = matched.second.back(); - auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) - ? cudf::nullable_join::YES - : cudf::nullable_join::NO; - - // For `inner_join`, we can freely choose either the `left` or `right` table to use for - // building/probing the hash map. Because building is typically more expensive than probing, we - // build the hash map from the smaller table. - if (right.num_rows() > left.num_rows()) { - cudf::hash_join hj_obj(left, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); - auto [right_result, left_result] = hj_obj.inner_join(right, std::nullopt, stream, mr); - return std::pair(std::move(left_result), std::move(right_result)); - } else { - cudf::hash_join hj_obj(right, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); + auto join_tables = [&](table_view const& left, table_view const& right) { + auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) + ? cudf::nullable_join::YES + : cudf::nullable_join::NO; + if (right.num_rows() > left.num_rows()) { + cudf::hash_join::impl_type hj_obj(left, + has_nulls == cudf::nullable_join::YES, + compare_nulls, + CUCO_DESIRED_LOAD_FACTOR, + stream, + cudf::get_current_device_resource_ref()); + auto [right_result, left_result] = hj_obj.inner_join(right, std::nullopt, stream, mr); + return std::pair(std::move(left_result), std::move(right_result)); + } + cudf::hash_join::impl_type hj_obj(right, + has_nulls == cudf::nullable_join::YES, + compare_nulls, + CUCO_DESIRED_LOAD_FACTOR, + stream, + cudf::get_current_device_resource_ref()); return hj_obj.inner_join(left, std::nullopt, stream, mr); - } + }; + + if (!has_dictionary_columns(left_input)) { return join_tables(left_input, right_input); } + + auto matched = cudf::dictionary::detail::match_dictionaries( + {left_input, right_input}, stream, cudf::get_current_device_resource_ref()); + return join_tables(matched.second.front(), matched.second.back()); } std::pair>, @@ -67,21 +81,24 @@ left_join(table_view const& left_input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - // Make sure any dictionary columns have matched key sets. - // This will return any new dictionary columns created as well as updated table_views. - auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, // these should match - stream, - cudf::get_current_device_resource_ref()); // temporary objects returned - // now rebuild the table views with the updated ones - table_view const left = matched.second.front(); - table_view const right = matched.second.back(); - auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) + auto join_tables = [&](table_view const& left, table_view const& right) { + auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) ? cudf::nullable_join::YES : cudf::nullable_join::NO; + cudf::hash_join::impl_type hj_obj(right, + has_nulls == cudf::nullable_join::YES, + compare_nulls, + CUCO_DESIRED_LOAD_FACTOR, + stream, + cudf::get_current_device_resource_ref()); + return hj_obj.left_join(left, std::nullopt, stream, mr); + }; - cudf::hash_join hj_obj(right, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); - return hj_obj.left_join(left, std::nullopt, stream, mr); + if (!has_dictionary_columns(left_input)) { return join_tables(left_input, right_input); } + + auto matched = cudf::dictionary::detail::match_dictionaries( + {left_input, right_input}, stream, cudf::get_current_device_resource_ref()); + return join_tables(matched.second.front(), matched.second.back()); } std::pair>, @@ -92,21 +109,24 @@ full_join(table_view const& left_input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - // Make sure any dictionary columns have matched key sets. - // This will return any new dictionary columns created as well as updated table_views. - auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, // these should match - stream, - cudf::get_current_device_resource_ref()); // temporary objects returned - // now rebuild the table views with the updated ones - table_view const left = matched.second.front(); - table_view const right = matched.second.back(); - auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) + auto join_tables = [&](table_view const& left, table_view const& right) { + auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) ? cudf::nullable_join::YES : cudf::nullable_join::NO; + cudf::hash_join::impl_type hj_obj(right, + has_nulls == cudf::nullable_join::YES, + compare_nulls, + CUCO_DESIRED_LOAD_FACTOR, + stream, + cudf::get_current_device_resource_ref()); + return hj_obj.full_join(left, std::nullopt, stream, mr); + }; + + if (!has_dictionary_columns(left_input)) { return join_tables(left_input, right_input); } - cudf::hash_join hj_obj(right, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); - return hj_obj.full_join(left, std::nullopt, stream, mr); + auto matched = cudf::dictionary::detail::match_dictionaries( + {left_input, right_input}, stream, cudf::get_current_device_resource_ref()); + return join_tables(matched.second.front(), matched.second.back()); } } // namespace detail From 0766c6f11c07716df3dee58b2f8b412dd225e3a8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 17:24:41 +0000 Subject: [PATCH 02/12] Remove HashCSR inline qualifiers --- cpp/src/join/hash_join/dispatch.cuh | 6 ++--- cpp/src/join/hash_join/hash_csr.cuh | 12 ++++----- cpp/src/join/hash_join/hash_csr_kernels.cuh | 30 +++++++++++---------- 3 files changed, 24 insertions(+), 24 deletions(-) diff --git a/cpp/src/join/hash_join/dispatch.cuh b/cpp/src/join/hash_join/dispatch.cuh index 9a5006d6d286..25d93c002da0 100644 --- a/cpp/src/join/hash_join/dispatch.cuh +++ b/cpp/src/join/hash_join/dispatch.cuh @@ -81,8 +81,7 @@ class hash_csr_pair_equal { { } - __device__ __forceinline__ bool operator()(probe_key_type const& lhs, - probe_key_type const& rhs) const noexcept + __device__ bool operator()(probe_key_type const& lhs, probe_key_type const& rhs) const noexcept { using detail::row::lhs_index_type; using detail::row::rhs_index_type; @@ -101,8 +100,7 @@ class primitive_hash_csr_pair_equal { { } - __device__ __forceinline__ bool operator()(probe_key_type const& lhs, - probe_key_type const& rhs) const noexcept + __device__ bool operator()(probe_key_type const& lhs, probe_key_type const& rhs) const noexcept { return _check_row_equality(lhs.second, rhs.second); } diff --git a/cpp/src/join/hash_join/hash_csr.cuh b/cpp/src/join/hash_join/hash_csr.cuh index 4a68b754f6cd..47b6684941e6 100644 --- a/cpp/src/join/hash_join/hash_csr.cuh +++ b/cpp/src/join/hash_join/hash_csr.cuh @@ -18,32 +18,32 @@ constexpr size_type hash_csr_empty_slot = size_type{-1}; constexpr std::uint64_t hash_csr_empty_entry = std::numeric_limits::max(); constexpr std::uint64_t hash_csr_empty_build_position = std::numeric_limits::max(); -__device__ inline std::uint64_t pack_hash_csr_entry(hash_value_type hash, size_type row) +static __device__ std::uint64_t pack_hash_csr_entry(hash_value_type hash, size_type row) { return (static_cast(hash) << 32) | static_cast(row); } -__device__ inline hash_value_type unpack_hash_csr_hash(std::uint64_t value) +static __device__ hash_value_type unpack_hash_csr_hash(std::uint64_t value) { return static_cast(value >> 32); } -__device__ inline size_type unpack_hash_csr_row(std::uint64_t value) +static __device__ size_type unpack_hash_csr_row(std::uint64_t value) { return static_cast(static_cast(value)); } -__device__ inline std::uint64_t pack_hash_csr_build_position(std::uint32_t slot, size_type rank) +static __device__ std::uint64_t pack_hash_csr_build_position(std::uint32_t slot, size_type rank) { return (static_cast(slot) << 32) | static_cast(rank); } -__device__ inline std::uint32_t unpack_hash_csr_build_slot(std::uint64_t value) +static __device__ std::uint32_t unpack_hash_csr_build_slot(std::uint64_t value) { return static_cast(value >> 32); } -__device__ inline size_type unpack_hash_csr_build_rank(std::uint64_t value) +static __device__ size_type unpack_hash_csr_build_rank(std::uint64_t value) { return static_cast(static_cast(value)); } diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index 67b7c4bad877..3d440d906706 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -135,11 +135,11 @@ void launch_hash_csr_build_count(size_type num_rows, CUDF_CUDA_TRY(cudaGetLastError()); } -inline void launch_hash_csr_build_fill(size_type num_rows, - std::uint64_t const* build_positions, - size_type const* cumulative_ends, - size_type* values, - rmm::cuda_stream_view stream) +[[maybe_unused]] static void launch_hash_csr_build_fill(size_type num_rows, + std::uint64_t const* build_positions, + size_type const* cumulative_ends, + size_type* values, + rmm::cuda_stream_view stream) { if (num_rows == 0) { return; } auto const config = grid_1d{num_rows, hash_csr_block_size}; @@ -208,7 +208,9 @@ void hash_csr_exclusive_scan(InputIterator input, stream.value())); } -inline void hash_csr_inclusive_sum(size_type* values, size_type size, rmm::cuda_stream_view stream) +[[maybe_unused]] static void hash_csr_inclusive_sum(size_type* values, + size_type size, + rmm::cuda_stream_view stream) { auto const mr = cudf::get_current_device_resource_ref(); std::size_t temp_storage_bytes{}; @@ -219,10 +221,10 @@ inline void hash_csr_inclusive_sum(size_type* values, size_type size, rmm::cuda_ temp_storage.data(), temp_storage_bytes, values, values, size, stream.value())); } -inline std::int64_t hash_csr_scan_counts(size_type const* counts, - size_type num_rows, - std::int64_t* offsets, - rmm::cuda_stream_view stream) +[[maybe_unused]] static std::int64_t hash_csr_scan_counts(size_type const* counts, + size_type num_rows, + std::int64_t* offsets, + rmm::cuda_stream_view stream) { auto const mr = cudf::get_current_device_resource_ref(); cudf::detail::device_scalar output_size(stream, mr); @@ -239,9 +241,9 @@ struct hash_csr_count_to_int64 { } }; -inline std::int64_t hash_csr_reduce_counts(size_type const* counts, - size_type num_rows, - rmm::cuda_stream_view stream) +[[maybe_unused]] static std::int64_t hash_csr_reduce_counts(size_type const* counts, + size_type num_rows, + rmm::cuda_stream_view stream) { if (num_rows == 0) { return 0; } auto const mr = cudf::get_current_device_resource_ref(); @@ -256,7 +258,7 @@ inline std::int64_t hash_csr_reduce_counts(size_type const* counts, return result.value(stream); } -__device__ inline size_type hash_csr_find_probe_row_in_range(std::int64_t const* offsets, +static __device__ size_type hash_csr_find_probe_row_in_range(std::int64_t const* offsets, size_type first_probe, size_type last_probe, std::int64_t output_index) From d9d4f94c983b6d9a4a0165bd5959a042cafb5fc4 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 17:28:50 +0000 Subject: [PATCH 03/12] Fix HashCSR memory resource ownership --- cpp/src/join/hash_join/hash_join.cu | 2 +- cpp/src/join/hash_join/hash_join_impl.cuh | 11 +++++++---- 2 files changed, 8 insertions(+), 5 deletions(-) diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index 027381ee817d..7e5bb9ba6ee3 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -103,7 +103,7 @@ hash_join::hash_join(cudf::table_view const& right, _right{right}, _preprocessed_right{cudf::detail::row::equality::preprocessed_table::create(_right, stream)}, _impl{std::make_unique( - hash_csr_capacity(right.num_rows(), load_factor), right.num_rows(), stream, mr)} + hash_csr_capacity(right.num_rows(), load_factor), right.num_rows(), stream, std::move(mr))} { CUDF_FUNC_RANGE(); CUDF_EXPECTS(0 != right.num_columns(), "Hash join right table is empty", std::invalid_argument); diff --git a/cpp/src/join/hash_join/hash_join_impl.cuh b/cpp/src/join/hash_join/hash_join_impl.cuh index 592b953d3211..96d0152eab34 100644 --- a/cpp/src/join/hash_join/hash_join_impl.cuh +++ b/cpp/src/join/hash_join/hash_join_impl.cuh @@ -12,6 +12,7 @@ #include #include +#include namespace cudf::detail { @@ -20,10 +21,11 @@ struct hash_join::impl { impl(std::uint32_t capacity, size_type rows, rmm::cuda_stream_view stream, - cuda::mr::any_resource const& mr) - : entries(capacity, stream, mr), - cumulative_ends(capacity, stream, mr), - values(rows, stream, mr), + cuda::mr::any_resource mr) + : _mr{std::move(mr)}, + entries(capacity, stream, _mr), + cumulative_ends(capacity, stream, _mr), + values(rows, stream, _mr), capacity{capacity} { } @@ -35,6 +37,7 @@ struct hash_join::impl { hash_csr_view csr_view() const { return {cumulative_ends.data(), values.data()}; } + cuda::mr::any_resource _mr; rmm::device_uvector entries; rmm::device_uvector cumulative_ends; rmm::device_uvector values; From 1d0cd216a1a4361bdc0b387fed842f0eebcf60e8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 17:38:53 +0000 Subject: [PATCH 04/12] Use one-phase CUB algorithms in HashCSR --- cpp/src/join/hash_join/hash_csr_kernels.cuh | 47 ++++++--------------- 1 file changed, 12 insertions(+), 35 deletions(-) diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index 3d440d906706..726fc24aab20 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -16,13 +16,13 @@ #include #include -#include -#include - #include #include +#include #include +#include #include +#include #include #include @@ -187,38 +187,19 @@ void hash_csr_exclusive_scan(InputIterator input, InitialValue initial_value, rmm::cuda_stream_view stream) { - auto const mr = cudf::get_current_device_resource_ref(); - std::size_t temp_storage_bytes{}; - CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan(nullptr, - temp_storage_bytes, - input, - output, - cuda::std::plus<>{}, - initial_value, - num_items, - stream.value())); - rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); - CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan(temp_storage.data(), - temp_storage_bytes, - input, - output, - cuda::std::plus<>{}, - initial_value, - num_items, - stream.value())); + auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, + cudf::get_current_device_resource_ref()}; + CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan( + input, output, cuda::std::plus<>{}, initial_value, num_items, env)); } [[maybe_unused]] static void hash_csr_inclusive_sum(size_type* values, size_type size, rmm::cuda_stream_view stream) { - auto const mr = cudf::get_current_device_resource_ref(); - std::size_t temp_storage_bytes{}; - CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum( - nullptr, temp_storage_bytes, values, values, size, stream.value())); - rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); - CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum( - temp_storage.data(), temp_storage_bytes, values, values, size, stream.value())); + auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, + cudf::get_current_device_resource_ref()}; + CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum(values, values, size, env)); } [[maybe_unused]] static std::int64_t hash_csr_scan_counts(size_type const* counts, @@ -249,12 +230,8 @@ struct hash_csr_count_to_int64 { auto const mr = cudf::get_current_device_resource_ref(); auto input = cuda::transform_iterator{counts, hash_csr_count_to_int64{}}; cudf::detail::device_scalar result(stream, mr); - std::size_t temp_storage_bytes{}; - CUDF_CUDA_TRY(cub::DeviceReduce::Sum( - nullptr, temp_storage_bytes, input, result.data(), num_rows, stream.value())); - rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); - CUDF_CUDA_TRY(cub::DeviceReduce::Sum( - temp_storage.data(), temp_storage_bytes, input, result.data(), num_rows, stream.value())); + auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, mr}; + CUDF_CUDA_TRY(cub::DeviceReduce::Sum(input, result.data(), num_rows, env)); return result.value(stream); } From 959b1883e062e0b8bad7039408fc420d7e64a099 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 17:45:11 +0000 Subject: [PATCH 05/12] Remove legacy hash join implementation --- cpp/CMakeLists.txt | 5 - cpp/src/join/hash_join/dispatch.cuh | 115 +------- cpp/src/join/hash_join/full_join_size_impl.cu | 100 ------- cpp/src/join/hash_join/hash_csr.cuh | 27 +- cpp/src/join/hash_join/hash_csr_kernels.cuh | 4 +- cpp/src/join/hash_join/hash_join.cu | 2 - cpp/src/join/hash_join/kernels_common.cuh | 24 -- cpp/src/join/hash_join/partitioned_count.cu | 25 -- .../hash_join/partitioned_count_kernels.cuh | 95 ------- .../hash_join/partitioned_count_kernels.hpp | 22 -- .../join/hash_join/partitioned_count_outer.cu | 25 -- .../join/hash_join/partitioned_retrieve.cu | 43 --- .../partitioned_retrieve_kernels.cuh | 258 ------------------ .../partitioned_retrieve_kernels.hpp | 41 --- .../hash_join/partitioned_retrieve_outer.cu | 43 --- cpp/src/join/hash_join/ref_types.cuh | 44 --- 16 files changed, 31 insertions(+), 842 deletions(-) delete mode 100644 cpp/src/join/hash_join/full_join_size_impl.cu delete mode 100644 cpp/src/join/hash_join/kernels_common.cuh delete mode 100644 cpp/src/join/hash_join/partitioned_count.cu delete mode 100644 cpp/src/join/hash_join/partitioned_count_kernels.cuh delete mode 100644 cpp/src/join/hash_join/partitioned_count_kernels.hpp delete mode 100644 cpp/src/join/hash_join/partitioned_count_outer.cu delete mode 100644 cpp/src/join/hash_join/partitioned_retrieve.cu delete mode 100644 cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh delete mode 100644 cpp/src/join/hash_join/partitioned_retrieve_kernels.hpp delete mode 100644 cpp/src/join/hash_join/partitioned_retrieve_outer.cu delete mode 100644 cpp/src/join/hash_join/ref_types.cuh diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 605e78eb0c14..47d3f05388db 100644 --- a/cpp/CMakeLists.txt +++ b/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 diff --git a/cpp/src/join/hash_join/dispatch.cuh b/cpp/src/join/hash_join/dispatch.cuh index 25d93c002da0..abe3d5b829e5 100644 --- a/cpp/src/join/hash_join/dispatch.cuh +++ b/cpp/src/join/hash_join/dispatch.cuh @@ -6,103 +6,47 @@ #include "common.cuh" -#include #include #include #include -#include - #include #include namespace cudf::detail { -/** - * @brief Equality comparator for cuco hash table probing with row-level equality. - */ -template -class pair_equal { - public: - pair_equal(Equal check_row_equality) : _check_row_equality{std::move(check_row_equality)} {} - - __device__ __forceinline__ bool operator()( - cuco::pair const& lhs, - cuco::pair const& rhs) const noexcept - { - using detail::row::lhs_index_type; - using detail::row::rhs_index_type; - - return lhs.first == rhs.first and - _check_row_equality(lhs_index_type{lhs.second}, rhs_index_type{rhs.second}); - } - - private: - 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 const& slot) const - { - return slot.second; - } -}; - -/** - * @brief Equality comparator for cuco hash table probing with primitive row equality. - */ -class primitive_pair_equal { - public: - primitive_pair_equal(cudf::detail::row::primitive::row_equality_comparator check_row_equality) - : _check_row_equality{std::move(check_row_equality)} - { - } - - __device__ __forceinline__ bool operator()( - cuco::pair const& lhs, - cuco::pair const& rhs) const noexcept - { - return lhs.first == rhs.first and _check_row_equality(lhs.second, rhs.second); - } - - private: - cudf::detail::row::primitive::row_equality_comparator _check_row_equality; -}; - template -class hash_csr_pair_equal { +class hash_csr_equal { public: - explicit hash_csr_pair_equal(Equal check_row_equality) + explicit hash_csr_equal(Equal check_row_equality) : _check_row_equality{std::move(check_row_equality)} { } - __device__ bool operator()(probe_key_type const& lhs, probe_key_type const& rhs) const noexcept + __device__ bool operator()(hash_csr_key_type const& lhs, + hash_csr_key_type const& rhs) const noexcept { using detail::row::lhs_index_type; using detail::row::rhs_index_type; - return _check_row_equality(lhs_index_type{lhs.second}, rhs_index_type{rhs.second}); + return _check_row_equality(lhs_index_type{lhs.row}, rhs_index_type{rhs.row}); } private: Equal _check_row_equality; }; -class primitive_hash_csr_pair_equal { +class primitive_hash_csr_equal { public: - explicit primitive_hash_csr_pair_equal( + explicit primitive_hash_csr_equal( cudf::detail::row::primitive::row_equality_comparator check_row_equality) : _check_row_equality{std::move(check_row_equality)} { } - __device__ bool operator()(probe_key_type const& lhs, probe_key_type const& rhs) const noexcept + __device__ bool operator()(hash_csr_key_type const& lhs, + hash_csr_key_type const& rhs) const noexcept { - return _check_row_equality(lhs.second, rhs.second); + return _check_row_equality(lhs.row, rhs.row); } private: @@ -125,40 +69,7 @@ decltype(auto) dispatch_hash_csr_comparator( auto const d_hasher = cudf::detail::row::primitive::row_hasher{left_nulls, preprocessed_left}; auto const d_equal = cudf::detail::row::primitive::row_equality_comparator{ left_nulls, preprocessed_left, preprocessed_right, compare_nulls}; - return std::forward(fn)(primitive_hash_csr_pair_equal{d_equal}, d_hasher); - } - - auto const d_hasher = - cudf::detail::row::hash::row_hasher{preprocessed_left}.device_hasher(left_nulls); - auto const row_comparator = - cudf::detail::row::equality::two_table_comparator{preprocessed_left, preprocessed_right}; - - if (cudf::detail::has_nested_columns(left_table)) { - auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); - return std::forward(fn)(hash_csr_pair_equal{d_equal}, d_hasher); - } - - auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); - return std::forward(fn)(hash_csr_pair_equal{d_equal}, d_hasher); -} - -template -decltype(auto) dispatch_join_comparator( - table_view const& right_table, - table_view const& left_table, - std::shared_ptr const& preprocessed_right, - std::shared_ptr const& preprocessed_left, - bool has_nulls, - null_equality compare_nulls, - Fn&& fn) -{ - auto const left_nulls = cudf::nullate::DYNAMIC{has_nulls}; - - if (cudf::detail::is_primitive_row_op_compatible(right_table)) { - auto const d_hasher = cudf::detail::row::primitive::row_hasher{left_nulls, preprocessed_left}; - auto const d_equal = cudf::detail::row::primitive::row_equality_comparator{ - left_nulls, preprocessed_left, preprocessed_right, compare_nulls}; - return std::forward(fn)(primitive_pair_equal{d_equal}, d_hasher); + return std::forward(fn)(primitive_hash_csr_equal{d_equal}, d_hasher); } auto const d_hasher = @@ -168,11 +79,11 @@ decltype(auto) dispatch_join_comparator( if (cudf::detail::has_nested_columns(left_table)) { auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); - return std::forward(fn)(pair_equal{d_equal}, d_hasher); + return std::forward(fn)(hash_csr_equal{d_equal}, d_hasher); } auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); - return std::forward(fn)(pair_equal{d_equal}, d_hasher); + return std::forward(fn)(hash_csr_equal{d_equal}, d_hasher); } } // namespace cudf::detail diff --git a/cpp/src/join/hash_join/full_join_size_impl.cu b/cpp/src/join/hash_join/full_join_size_impl.cu deleted file mode 100644 index 13bb1393f64e..000000000000 --- a/cpp/src/join/hash_join/full_join_size_impl.cu +++ /dev/null @@ -1,100 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#if 0 // Replaced by the HashCSR full-join size path in size_impl.cuh. -#include "retrieve_impl.cuh" - -#include -#include - -#include -#include - -#include -#include -#include -#include - -#include - -namespace cudf::detail { - -namespace { -std::size_t compute_left_join_complement_size(cudf::device_span right_indices, - size_type left_table_row_count, - size_type right_table_row_count, - rmm::cuda_stream_view stream) -{ - if (left_table_row_count == 0) { return right_table_row_count; } - - auto invalid_index_map = - std::make_unique>(right_table_row_count, stream); - thrust::uninitialized_fill( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - invalid_index_map->begin(), - invalid_index_map->end(), - int32_t{1}); - - valid_range valid(0, right_table_row_count); - - thrust::scatter_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::make_constant_iterator(0), - cuda::make_constant_iterator(0) + right_indices.size(), - right_indices.begin(), - right_indices.begin(), - invalid_index_map->begin(), - valid); - - return cudf::detail::count_if( - invalid_index_map->begin(), invalid_index_map->end(), cuda::std::identity{}, stream); -} -} // namespace - -std::size_t get_full_join_size( - cudf::table_view const& right_table, - cudf::table_view const& left_table, - std::shared_ptr const& preprocessed_right, - std::shared_ptr const& preprocessed_left, - cudf::detail::hash_table_t const& hash_table, - bool has_nulls, - null_equality compare_nulls, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - std::size_t join_size = compute_join_output_size(right_table, - left_table, - preprocessed_right, - preprocessed_left, - hash_table, - has_nulls, - compare_nulls, - stream); - - if (join_size == 0) { return join_size; } - - auto right_indices = std::make_unique>(join_size, stream, mr); - - auto const out_build_begin = - cuda::make_transform_output_iterator(right_indices->begin(), output_fn{}); - - retrieve_left_join_build_indices(right_table, - left_table, - preprocessed_right, - preprocessed_left, - hash_table, - has_nulls, - compare_nulls, - out_build_begin, - stream); - - auto const left_table_row_count = left_table.num_rows(); - auto const right_table_row_count = right_table.num_rows(); - - return join_size + compute_left_join_complement_size( - *right_indices, left_table_row_count, right_table_row_count, stream); -} - -} // namespace cudf::detail -#endif diff --git a/cpp/src/join/hash_join/hash_csr.cuh b/cpp/src/join/hash_join/hash_csr.cuh index 47b6684941e6..25c5b96c31ec 100644 --- a/cpp/src/join/hash_join/hash_csr.cuh +++ b/cpp/src/join/hash_join/hash_csr.cuh @@ -5,8 +5,7 @@ #pragma once -#include "kernels_common.cuh" - +#include #include #include @@ -14,6 +13,11 @@ namespace cudf::detail { +struct hash_csr_key_type { + hash_value_type hash; + size_type row; +}; + constexpr size_type hash_csr_empty_slot = size_type{-1}; constexpr std::uint64_t hash_csr_empty_entry = std::numeric_limits::max(); constexpr std::uint64_t hash_csr_empty_build_position = std::numeric_limits::max(); @@ -54,17 +58,17 @@ struct hash_csr_map_view { std::uint32_t mask; template - __device__ std::uint32_t find_or_insert(probe_key_type key, Equal equal) const + __device__ std::uint32_t find_or_insert(hash_csr_key_type key, Equal equal) const { - auto const desired = pack_hash_csr_entry(key.first, key.second); + auto const desired = pack_hash_csr_entry(key.hash, key.row); for (std::uint32_t step = 0; step < capacity; ++step) { - auto const slot = (static_cast(key.first) + step) & mask; + auto const slot = (static_cast(key.hash) + step) & mask; auto const old = atomicCAS(reinterpret_cast(entries + slot), static_cast(hash_csr_empty_entry), static_cast(desired)); if (old == hash_csr_empty_entry) { return slot; } - if (unpack_hash_csr_hash(old) == key.first && - equal(key, probe_key_type{unpack_hash_csr_hash(old), unpack_hash_csr_row(old)})) { + if (unpack_hash_csr_hash(old) == key.hash && + equal(key, hash_csr_key_type{unpack_hash_csr_hash(old), unpack_hash_csr_row(old)})) { return slot; } } @@ -72,14 +76,15 @@ struct hash_csr_map_view { } template - __device__ std::uint32_t find(probe_key_type key, Equal equal) const + __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(key.first) + step) & mask; + auto const slot = (static_cast(key.hash) + step) & mask; auto const current = entries[slot]; if (current == hash_csr_empty_entry) { return capacity; } - if (unpack_hash_csr_hash(current) == key.first && - equal(key, probe_key_type{unpack_hash_csr_hash(current), unpack_hash_csr_row(current)})) { + if (unpack_hash_csr_hash(current) == key.hash && + equal(key, + hash_csr_key_type{unpack_hash_csr_hash(current), unpack_hash_csr_row(current)})) { return slot; } } diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index 726fc24aab20..cf38f11f617a 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -52,7 +52,7 @@ CUDF_KERNEL void hash_csr_build_count_kernel(size_type num_rows, continue; } - auto const slot = map.find_or_insert(probe_key_type{hasher(index), index}, equal); + auto const slot = map.find_or_insert(hash_csr_key_type{hasher(index), index}, equal); if (slot == map.capacity) { build_positions[index] = hash_csr_empty_build_position; continue; @@ -96,7 +96,7 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, auto const index = static_cast(row); auto slot = map.capacity; if (valid_rows == nullptr || cudf::bit_is_set(valid_rows, index)) { - slot = map.find(probe_key_type{hasher(index), index}, equal); + slot = map.find(hash_csr_key_type{hasher(index), index}, equal); } auto const found = slot != map.capacity; diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index 7e5bb9ba6ee3..2560bdc4264c 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -21,8 +21,6 @@ #include #include -#include - #include #include diff --git a/cpp/src/join/hash_join/kernels_common.cuh b/cpp/src/join/hash_join/kernels_common.cuh deleted file mode 100644 index cd74a5bde4e1..000000000000 --- a/cpp/src/join/hash_join/kernels_common.cuh +++ /dev/null @@ -1,24 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -// Custom hash-join probe kernels that give cudf direct control over kernel launches. -// Uses the cuco ref type for hash-table access (storage, probing scheme, predicate). - -#pragma once - -#include "join/join_common_utils.hpp" - -#include -#include -#include - -#include - -namespace cudf::detail { - -/// The probe key type stored in the hash table: {hash_value, row_index}. -using probe_key_type = cuco::pair; - -} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/partitioned_count.cu b/cpp/src/join/hash_join/partitioned_count.cu deleted file mode 100644 index 345f9f9a7053..000000000000 --- a/cpp/src/join/hash_join/partitioned_count.cu +++ /dev/null @@ -1,25 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#if 0 // Replaced by the HashCSR match-context count path. -#include "partitioned_count_kernels.cuh" -#include "ref_types.cuh" - -namespace cudf::detail { - -template void launch_partitioned_count(probe_key_type const*, - thread_index_type, - size_type*, - primitive_count_ref_t, - rmm::cuda_stream_view); - -template void launch_partitioned_count( - probe_key_type const*, thread_index_type, size_type*, nested_count_ref_t, rmm::cuda_stream_view); - -template void launch_partitioned_count( - probe_key_type const*, thread_index_type, size_type*, flat_count_ref_t, rmm::cuda_stream_view); - -} // namespace cudf::detail -#endif diff --git a/cpp/src/join/hash_join/partitioned_count_kernels.cuh b/cpp/src/join/hash_join/partitioned_count_kernels.cuh deleted file mode 100644 index cf21086b68e9..000000000000 --- a/cpp/src/join/hash_join/partitioned_count_kernels.cuh +++ /dev/null @@ -1,95 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include "kernels_common.cuh" - -#include - -#include - -#include -#include - -namespace cudf::detail { - -/** - * @brief Count matching build-side rows for each probe key. - * - * Each probing tile (@p cg_size threads) calls `ref.count()` for one probe key - * and reduces the per-lane counts across the tile with a warp reduce. The result - * is written to @p output by a single elected thread via `invoke_one`. If - * @p IsOuter is true, keys with zero matches are recorded as 1 so every probe - * row contributes at least one output row in the subsequent retrieve pass. - * - * This is the first phase of the two-phase partitioned join: count then retrieve. - * The output array is consumed by `launch_partitioned_retrieve` to pre-allocate - * the output index buffers. - * - * @tparam IsOuter If true, zero-match keys produce a count of 1 - * @tparam Ref cuco open-addressing reference type (carries hash, equality, storage) - * @param keys Packed probe keys: `.first` = hash, `.second` = probe row index - * @param n Number of probe keys - * @param output Per-key match count output (one entry per probe key) - * @param ref cuco hash-table reference for counting - */ -template -CUDF_KERNEL void __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) - partitioned_count_kernel(probe_key_type const* __restrict__ keys, - thread_index_type n, - size_type* __restrict__ output, - Ref ref) -{ - auto constexpr cg_size = DEFAULT_JOIN_CG_SIZE; - - auto idx = grid_1d::global_thread_id() / cg_size; - auto const stride = grid_1d::grid_stride() / cg_size; - - while (idx < n) { - auto const key = keys[idx]; - if constexpr (cg_size == 1) { - auto const match_count = ref.count(key); - if constexpr (IsOuter) { - output[idx] = (match_count == 0) ? size_type{1} : match_count; - } else { - output[idx] = match_count; - } - } else { - auto const tile = - cooperative_groups::tiled_partition(cooperative_groups::this_thread_block()); - auto const temp_count = static_cast(ref.count(tile, key)); - auto const match_count = - cooperative_groups::reduce(tile, temp_count, cooperative_groups::plus()); - cooperative_groups::invoke_one(tile, [&]() { - if constexpr (IsOuter) { - output[idx] = (match_count == 0) ? size_type{1} : match_count; - } else { - output[idx] = match_count; - } - }); - } - idx += stride; - } -} - -template -void launch_partitioned_count(probe_key_type const* keys, - thread_index_type n, - size_type* output, - Ref ref, - rmm::cuda_stream_view stream) -{ - if (n == 0) { return; } - - auto const config = - grid_1d{static_cast(n * DEFAULT_JOIN_CG_SIZE), DEFAULT_JOIN_BLOCK_SIZE}; - - partitioned_count_kernel - <<>>(keys, n, output, ref); - CUDF_CUDA_TRY(cudaGetLastError()); -} - -} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/partitioned_count_kernels.hpp b/cpp/src/join/hash_join/partitioned_count_kernels.hpp deleted file mode 100644 index d6bb30b7f00f..000000000000 --- a/cpp/src/join/hash_join/partitioned_count_kernels.hpp +++ /dev/null @@ -1,22 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include "kernels_common.cuh" - -#include - -namespace cudf::detail { - -/// Launch the partitioned_count kernel. -template -void launch_partitioned_count(probe_key_type const* keys, - thread_index_type n, - size_type* output, - Ref ref, - rmm::cuda_stream_view stream); - -} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/partitioned_count_outer.cu b/cpp/src/join/hash_join/partitioned_count_outer.cu deleted file mode 100644 index d12c775e2221..000000000000 --- a/cpp/src/join/hash_join/partitioned_count_outer.cu +++ /dev/null @@ -1,25 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#if 0 // Replaced by the shared HashCSR outer count kernel. -#include "partitioned_count_kernels.cuh" -#include "ref_types.cuh" - -namespace cudf::detail { - -template void launch_partitioned_count(probe_key_type const*, - thread_index_type, - size_type*, - primitive_count_ref_t, - rmm::cuda_stream_view); - -template void launch_partitioned_count( - probe_key_type const*, thread_index_type, size_type*, nested_count_ref_t, rmm::cuda_stream_view); - -template void launch_partitioned_count( - probe_key_type const*, thread_index_type, size_type*, flat_count_ref_t, rmm::cuda_stream_view); - -} // namespace cudf::detail -#endif diff --git a/cpp/src/join/hash_join/partitioned_retrieve.cu b/cpp/src/join/hash_join/partitioned_retrieve.cu deleted file mode 100644 index 29ba4bc9e0b0..000000000000 --- a/cpp/src/join/hash_join/partitioned_retrieve.cu +++ /dev/null @@ -1,43 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#if 0 // Replaced by the HashCSR partitioned pull path. -#include "partitioned_retrieve_kernels.cuh" -#include "ref_types.cuh" - -namespace cudf::detail { - -template std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const*, - thread_index_type, - size_type const*, - primitive_count_ref_t, - size_type, - rmm::cuda_stream_view, - rmm::device_async_resource_ref); - -template std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const*, - thread_index_type, - size_type const*, - nested_count_ref_t, - size_type, - rmm::cuda_stream_view, - rmm::device_async_resource_ref); - -template std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const*, - thread_index_type, - size_type const*, - flat_count_ref_t, - size_type, - rmm::cuda_stream_view, - rmm::device_async_resource_ref); - -} // namespace cudf::detail -#endif diff --git a/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh b/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh deleted file mode 100644 index f86e990bfe9a..000000000000 --- a/cpp/src/join/hash_join/partitioned_retrieve_kernels.cuh +++ /dev/null @@ -1,258 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include "kernels_common.cuh" - -#include -#include -#include -#include - -#include -#include -#include - -#include -#include -#include -#include -#include -#include - -namespace cudf::detail { - -/** - * @brief Count the number of set bits below a given position in a bitmask. - */ -__device__ __forceinline__ int count_lower_set_bits(unsigned int mask, int pos) -{ - return cuda::std::popcount(mask & ((1u << pos) - 1)); -} - -/** - * @brief Retrieve matching build-side rows for each probe key. - * - * Each probing tile (@p cg_size threads) walks the hash table for one probe key, - * collecting matches via warp ballot. Matches are staged in a per-flushing-tile (warp) - * shared-memory buffer instead of being written directly to global memory. When the buffer - * nears capacity, the flushing tile claims a contiguous range in the global output arrays - * via a single atomic and flushes with coalesced writes, amortising atomic overhead across - * many matches. If @p IsOuter is true, probe rows with no matches emit a - * `(left_index, JoinNoMatch)` pair. - * - * @tparam IsOuter If true, unmatched probe rows emit a null-padded output row - * @tparam Ref cuco open-addressing reference type (carries hash, equality, storage) - * @param keys Packed left keys: `.first` = hash, `.second` = left row index - * @param n Number of probe keys - * @param left_offset Added to each probe row index to produce an absolute left index - * @param left_output Output buffer for left (probe-side) row indices - * @param right_output Output buffer for right (build-side) row indices - * @param output_counter Global atomic counter tracking total pairs written so far - * @param ref cuco hash-table reference for probing - */ -template -CUDF_KERNEL void __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) - partitioned_retrieve_kernel(probe_key_type const* __restrict__ keys, - thread_index_type n, - size_type left_offset, - size_type* __restrict__ left_output, - size_type* __restrict__ right_output, - size_type* __restrict__ output_counter, - Ref ref) -{ - namespace cg = cooperative_groups; - - auto constexpr cg_size = Ref::cg_size; - auto constexpr bucket_size = Ref::bucket_size; - auto constexpr flushing_tile_size = 32; // full warp for coalesced flushes - static_assert(flushing_tile_size >= cg_size); - static_assert(flushing_tile_size % cg_size == 0, - "Every probing tile must sit inside a single flushing tile"); - static_assert(DEFAULT_JOIN_BLOCK_SIZE % flushing_tile_size == 0); - - auto constexpr num_flushing_tiles = DEFAULT_JOIN_BLOCK_SIZE / flushing_tile_size; - auto constexpr tiles_in_block = DEFAULT_JOIN_BLOCK_SIZE / cg_size; - auto constexpr max_matches_per_step = flushing_tile_size * bucket_size; - // buffer_size leaves headroom so one full probing step can't overflow. - auto constexpr buffer_size = max_matches_per_step + flushing_tile_size; - - using index_pair = cuco::pair; - __shared__ index_pair buffers[num_flushing_tiles][buffer_size]; - __shared__ cuda::std::int32_t counters[num_flushing_tiles]; - - auto const block = cg::this_thread_block(); - auto const flushing_tile = cg::tiled_partition(block); - auto const probing_tile = cg::tiled_partition(block); - auto const flushing_tile_id = flushing_tile.meta_group_rank(); - auto const empty_sentinel = ref.empty_key_sentinel(); - auto const key_equal = ref.key_eq(); - - if (flushing_tile.thread_rank() == 0) { counters[flushing_tile_id] = 0; } - flushing_tile.sync(); - - auto atomic_counter = cuda::atomic_ref{*output_counter}; - - auto flush_buffers = [&](auto const& tile) { - auto const count = counters[flushing_tile_id]; - auto const offset = cg::invoke_one_broadcast(tile, [&]() { - return atomic_counter.fetch_add(static_cast(count), cuda::memory_order_relaxed); - }); - auto const rank = tile.thread_rank(); - for (int i = rank; i < count; i += tile.size()) { - left_output[offset + i] = buffers[flushing_tile_id][i].first; - right_output[offset + i] = buffers[flushing_tile_id][i].second; - } - }; - - auto const grid_stride_tiles = static_cast(gridDim.x) * tiles_in_block; - auto idx = - static_cast(blockIdx.x) * tiles_in_block + probing_tile.meta_group_rank(); - - while (flushing_tile.any(idx < n)) { - bool const active = idx < n; - auto const active_flushing_tile = - cg::binary_partition(flushing_tile, active); - - if (active) { - auto const probe_key = keys[idx]; - auto const left_index = probe_key.second + left_offset; - - auto probing_iter = ref.probing_scheme().template make_iterator( - probing_tile, probe_key, ref.storage_ref().extent()); - auto const init_probing_idx = *probing_iter; - - bool running = true; - [[maybe_unused]] bool found_match = false; - - while (active_flushing_tile.any(running)) { - if (running) { - auto const bucket_slots = ref.storage_ref()[*probing_iter]; - - bool equals[bucket_size]; - for (int i = 0; i < bucket_size; ++i) { - equals[i] = false; - if (running) { - if (bucket_slots[i] == empty_sentinel) { - running = false; - } else if (key_equal(probe_key, bucket_slots[i])) { - equals[i] = true; - } - } - } - - probing_tile.sync(); - running = probing_tile.all(running); - - cuda::std::int32_t exists[bucket_size]; - cuda::std::int32_t num_matches[bucket_size]; - cuda::std::int32_t total_matches = 0; - for (int i = 0; i < bucket_size; ++i) { - exists[i] = probing_tile.ballot(equals[i]); - num_matches[i] = cuda::std::popcount(static_cast(exists[i])); - total_matches += num_matches[i]; - } - - auto const lane_id = probing_tile.thread_rank(); - - if (total_matches > 0) { - if constexpr (IsOuter) { found_match = true; } - - cuda::std::int32_t output_idx = 0; - if (lane_id == 0) { - auto shared_ref = cuda::atomic_ref{ - counters[flushing_tile_id]}; - output_idx = shared_ref.fetch_add(total_matches, cuda::memory_order_relaxed); - } - output_idx = probing_tile.shfl(output_idx, 0); - - cuda::std::int32_t matches_offset = 0; - for (int i = 0; i < bucket_size; ++i) { - if (equals[i]) { - auto const lane_offset = count_lower_set_bits(exists[i], lane_id); - buffers[flushing_tile_id][output_idx + matches_offset + lane_offset] = { - left_index, bucket_slots[i].second}; - } - matches_offset += num_matches[i]; - } - } - - if constexpr (IsOuter) { - if (!running && !found_match && lane_id == 0) { - auto shared_ref = cuda::atomic_ref{ - counters[flushing_tile_id]}; - auto const output_idx = shared_ref.fetch_add(1, cuda::memory_order_relaxed); - buffers[flushing_tile_id][output_idx] = {left_index, cudf::JoinNoMatch}; - } - } - } // if running - - active_flushing_tile.sync(); - if (counters[flushing_tile_id] > (buffer_size - max_matches_per_step)) { - flush_buffers(active_flushing_tile); - active_flushing_tile.sync(); - if (active_flushing_tile.thread_rank() == 0) { counters[flushing_tile_id] = 0; } - active_flushing_tile.sync(); - } - - ++probing_iter; - if (*probing_iter == init_probing_idx) { running = false; } - } // while running - } // if active - - idx += grid_stride_tiles; - } // while idx < n - - flushing_tile.sync(); - if (counters[flushing_tile_id] > 0) { flush_buffers(flushing_tile); } -} - -template -std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const* keys, - thread_index_type n, - size_type const* match_counts, - Ref ref, - size_type left_offset, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - if (n == 0) { - return std::pair(std::make_unique>(0, stream, mr), - std::make_unique>(0, stream, mr)); - } - - // Shared-memory buffered retrieve only needs the total output size, not - // per-row offsets. A reduce is cheaper than an exclusive_scan. - auto const total_output = - thrust::reduce(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - match_counts, - match_counts + n, - size_type{0}); - - if (total_output == 0) { - return std::pair(std::make_unique>(0, stream, mr), - std::make_unique>(0, stream, mr)); - } - - auto left_indices = std::make_unique>(total_output, stream, mr); - auto right_indices = std::make_unique>(total_output, stream, mr); - - // Global atomic counter claimed in bulk by each flushing-tile buffer flush. - cudf::detail::device_scalar output_counter(size_type{0}, stream); - - auto constexpr tiles_in_block = DEFAULT_JOIN_BLOCK_SIZE / Ref::cg_size; - auto const num_blocks = static_cast((n + tiles_in_block - 1) / tiles_in_block); - - partitioned_retrieve_kernel<<>>( - keys, n, left_offset, left_indices->data(), right_indices->data(), output_counter.data(), ref); - CUDF_CUDA_TRY(cudaGetLastError()); - - return std::pair(std::move(left_indices), std::move(right_indices)); -} - -} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/partitioned_retrieve_kernels.hpp b/cpp/src/join/hash_join/partitioned_retrieve_kernels.hpp deleted file mode 100644 index edd8816924ad..000000000000 --- a/cpp/src/join/hash_join/partitioned_retrieve_kernels.hpp +++ /dev/null @@ -1,41 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -#pragma once - -#include "kernels_common.cuh" - -#include - -#include -#include - -#include -#include - -namespace cudf::detail { - -/** - * @brief Probes the hash table for each key and writes matching index pairs. - * - * Reduces match_counts to derive the total output size, allocates output buffers, - * and launches the retrieve kernel. `left_offset` is added to each stored probe-row index when - * writing to `left_indices`, so callers can produce indices in the full probe - * table's coordinate space directly from a slice-local `keys` array. - * - * @return A pair of device vectors [left_indices, right_indices]. - */ -template -std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const* keys, - thread_index_type n, - size_type const* match_counts, - Ref ref, - size_type left_offset, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); - -} // namespace cudf::detail diff --git a/cpp/src/join/hash_join/partitioned_retrieve_outer.cu b/cpp/src/join/hash_join/partitioned_retrieve_outer.cu deleted file mode 100644 index 1f6a519d78ba..000000000000 --- a/cpp/src/join/hash_join/partitioned_retrieve_outer.cu +++ /dev/null @@ -1,43 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. - * SPDX-License-Identifier: Apache-2.0 - */ - -#if 0 // Replaced by the shared HashCSR outer pull kernel. -#include "partitioned_retrieve_kernels.cuh" -#include "ref_types.cuh" - -namespace cudf::detail { - -template std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const*, - thread_index_type, - size_type const*, - primitive_count_ref_t, - size_type, - rmm::cuda_stream_view, - rmm::device_async_resource_ref); - -template std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const*, - thread_index_type, - size_type const*, - nested_count_ref_t, - size_type, - rmm::cuda_stream_view, - rmm::device_async_resource_ref); - -template std::pair>, - std::unique_ptr>> -launch_partitioned_retrieve(probe_key_type const*, - thread_index_type, - size_type const*, - flat_count_ref_t, - size_type, - rmm::cuda_stream_view, - rmm::device_async_resource_ref); - -} // namespace cudf::detail -#endif diff --git a/cpp/src/join/hash_join/ref_types.cuh b/cpp/src/join/hash_join/ref_types.cuh deleted file mode 100644 index a1511c56dcb1..000000000000 --- a/cpp/src/join/hash_join/ref_types.cuh +++ /dev/null @@ -1,44 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. - * SPDX-License-Identifier: Apache-2.0 - */ - -// Type aliases for the cuco hash table ref types and equality comparators -// used across hash join probe kernels. There are 3 dispatch paths: -// primitive, nested, non-nested. - -#pragma once - -#include "dispatch.cuh" -#include "hash_join_impl.cuh" - -#include - -#include - -namespace cudf::detail { - -// --- Equality types from the 3 dispatch paths --- - -using primitive_equality_t = primitive_pair_equal; - -using nested_equality_t = pair_equal>>; - -using flat_equality_t = pair_equal>>; - -// --- Count ref types (used by partitioned_count kernel) --- - -template -using count_ref_t = - decltype(std::declval() - .ref(cuco::op::count) - .rebind_key_eq(std::declval()) - .rebind_hash_function(std::declval().hash_function())); - -using primitive_count_ref_t = count_ref_t; -using nested_count_ref_t = count_ref_t; -using flat_count_ref_t = count_ref_t; - -} // namespace cudf::detail From e0b1d87d883210673e9ff0ba97e1321fb72053af Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 18:36:21 +0000 Subject: [PATCH 06/12] Rename HashCSR insert method --- cpp/src/join/hash_join/hash_csr.cuh | 2 +- cpp/src/join/hash_join/hash_csr_kernels.cuh | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/join/hash_join/hash_csr.cuh b/cpp/src/join/hash_join/hash_csr.cuh index 25c5b96c31ec..d0d95dae1580 100644 --- a/cpp/src/join/hash_join/hash_csr.cuh +++ b/cpp/src/join/hash_join/hash_csr.cuh @@ -58,7 +58,7 @@ struct hash_csr_map_view { std::uint32_t mask; template - __device__ std::uint32_t find_or_insert(hash_csr_key_type key, Equal equal) const + __device__ std::uint32_t insert(hash_csr_key_type key, Equal equal) const { auto const desired = pack_hash_csr_entry(key.hash, key.row); for (std::uint32_t step = 0; step < capacity; ++step) { diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index cf38f11f617a..5d2dd2c216cf 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -52,7 +52,7 @@ CUDF_KERNEL void hash_csr_build_count_kernel(size_type num_rows, continue; } - auto const slot = map.find_or_insert(hash_csr_key_type{hasher(index), index}, equal); + auto const slot = map.insert(hash_csr_key_type{hasher(index), index}, equal); if (slot == map.capacity) { build_positions[index] = hash_csr_empty_build_position; continue; From 164dcf0a3f379a816a9487fe8f9f331d490cad7c Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 18:37:43 +0000 Subject: [PATCH 07/12] Use cuDF warp size in HashCSR --- cpp/src/join/hash_join/hash_csr_kernels.cuh | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index 5d2dd2c216cf..332b5c8016b5 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 #include #include +#include #include #include #include @@ -30,9 +31,9 @@ namespace cudf::detail { -constexpr thread_index_type hash_csr_block_size = 256; -constexpr thread_index_type hash_csr_warp_size = 32; -constexpr thread_index_type hash_csr_warps_per_block = hash_csr_block_size / hash_csr_warp_size; +constexpr thread_index_type hash_csr_block_size = 256; +constexpr thread_index_type hash_csr_warps_per_block = + hash_csr_block_size / cudf::detail::warp_size; constexpr thread_index_type hash_csr_outputs_per_lane = 32; template @@ -264,8 +265,8 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, size_type* left_indices, size_type* right_indices) { - auto const lane_id = static_cast(threadIdx.x) % hash_csr_warp_size; - auto const warp_in_block = static_cast(threadIdx.x) / hash_csr_warp_size; + auto const lane_id = static_cast(threadIdx.x) % cudf::detail::warp_size; + auto const warp_in_block = static_cast(threadIdx.x) / cudf::detail::warp_size; auto const global_warp = static_cast(blockIdx.x) * hash_csr_warps_per_block + warp_in_block; auto const range_begin = outputs_per_warp * global_warp; @@ -282,7 +283,7 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, #pragma unroll for (thread_index_type item = 0; item < hash_csr_outputs_per_lane; ++item) { - auto const output_index = range_begin + lane_id + item * hash_csr_warp_size; + auto const output_index = range_begin + lane_id + item * cudf::detail::warp_size; if (output_index < range_end) { auto const probe_row = first_probe == last_probe @@ -316,7 +317,7 @@ void launch_hash_csr_retrieve(std::int64_t output_size, if (output_size == 0) { return; } auto const min_blocks = size_type{2} * cudf::detail::num_multiprocessors(); constexpr auto outputs_per_block = - hash_csr_warps_per_block * hash_csr_warp_size * hash_csr_outputs_per_lane; + hash_csr_warps_per_block * cudf::detail::warp_size * hash_csr_outputs_per_lane; auto const requested_blocks = (output_size + outputs_per_block - 1) / outputs_per_block; auto const num_blocks = static_cast(cuda::std::max(requested_blocks, min_blocks)); From 3ef3477a4893bbcfc20beecce097a079bbe4644b Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 18:40:47 +0000 Subject: [PATCH 08/12] Use atomic refs in HashCSR --- cpp/src/join/hash_join/hash_csr.cuh | 11 +++++++---- cpp/src/join/hash_join/hash_csr_kernels.cuh | 15 +++++++++++---- 2 files changed, 18 insertions(+), 8 deletions(-) diff --git a/cpp/src/join/hash_join/hash_csr.cuh b/cpp/src/join/hash_join/hash_csr.cuh index d0d95dae1580..4be73a5efe64 100644 --- a/cpp/src/join/hash_join/hash_csr.cuh +++ b/cpp/src/join/hash_join/hash_csr.cuh @@ -8,6 +8,8 @@ #include #include +#include + #include #include @@ -63,10 +65,11 @@ struct hash_csr_map_view { auto const desired = pack_hash_csr_entry(key.hash, key.row); for (std::uint32_t step = 0; step < capacity; ++step) { auto const slot = (static_cast(key.hash) + step) & mask; - auto const old = atomicCAS(reinterpret_cast(entries + slot), - static_cast(hash_csr_empty_entry), - static_cast(desired)); - if (old == hash_csr_empty_entry) { return slot; } + auto entry_ref = cuda::atomic_ref{entries[slot]}; + auto old = hash_csr_empty_entry; + if (entry_ref.compare_exchange_strong(old, desired, cuda::memory_order_relaxed)) { + return slot; + } if (unpack_hash_csr_hash(old) == key.hash && equal(key, hash_csr_key_type{unpack_hash_csr_hash(old), unpack_hash_csr_row(old)})) { return slot; diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index 332b5c8016b5..a671636b3eaa 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -58,7 +58,8 @@ CUDF_KERNEL void hash_csr_build_count_kernel(size_type num_rows, build_positions[index] = hash_csr_empty_build_position; continue; } - auto const rank = atomicAdd(slot_counts + slot, size_type{1}); + auto slot_count_ref = cuda::atomic_ref{slot_counts[slot]}; + auto const rank = slot_count_ref.fetch_add(size_type{1}, cuda::memory_order_relaxed); build_positions[index] = pack_hash_csr_build_position(slot, rank); } } @@ -109,9 +110,15 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, match_counts[index] = is_outer ? cuda::std::max(count, size_type{1}) : count; } - if (found && matched_slots != nullptr && - atomicCAS(matched_slots + slot, std::uint32_t{0}, std::uint32_t{1}) == 0) { - atomicAdd(matched_build_rows, static_cast(count)); + if (found && matched_slots != nullptr) { + auto matched_slot_ref = + cuda::atomic_ref{matched_slots[slot]}; + auto expected = std::uint32_t{0}; + if (matched_slot_ref.compare_exchange_strong( + expected, std::uint32_t{1}, cuda::memory_order_relaxed)) { + cuda::atomic_ref{*matched_build_rows} + .fetch_add(static_cast(count), cuda::memory_order_relaxed); + } } } } From 5c07e4411bdc73047fcfe6be683ddad8681fcde8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 18:42:40 +0000 Subject: [PATCH 09/12] Remove HashCSR scan wrappers --- cpp/src/join/hash_join/hash_csr_kernels.cuh | 26 +++------------------ cpp/src/join/hash_join/hash_join.cu | 5 +++- 2 files changed, 7 insertions(+), 24 deletions(-) diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index a671636b3eaa..3d753e8250e0 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -188,28 +188,6 @@ void launch_hash_csr_probe_count(size_type num_rows, CUDF_CUDA_TRY(cudaGetLastError()); } -template -void hash_csr_exclusive_scan(InputIterator input, - OutputIterator output, - size_type num_items, - InitialValue initial_value, - rmm::cuda_stream_view stream) -{ - auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, - cudf::get_current_device_resource_ref()}; - CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan( - input, output, cuda::std::plus<>{}, initial_value, num_items, env)); -} - -[[maybe_unused]] static void hash_csr_inclusive_sum(size_type* values, - size_type size, - rmm::cuda_stream_view stream) -{ - auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, - cudf::get_current_device_resource_ref()}; - CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum(values, values, size, env)); -} - [[maybe_unused]] static std::int64_t hash_csr_scan_counts(size_type const* counts, size_type num_rows, std::int64_t* offsets, @@ -219,7 +197,9 @@ void hash_csr_exclusive_scan(InputIterator input, cudf::detail::device_scalar output_size(stream, mr); auto output = cudf::detail::make_sizes_to_offsets_iterator( offsets, offsets + num_rows + 1, output_size.data()); - hash_csr_exclusive_scan(counts, output, num_rows + 1, std::int64_t{0}, stream); + auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, mr}; + CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan( + counts, output, cuda::std::plus<>{}, std::int64_t{0}, num_rows + 1, env)); return output_size.value(stream); } diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index 2560bdc4264c..905a38578503 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -132,7 +132,10 @@ hash_join::hash_join(cudf::table_view const& right, }; dispatch_hash_csr_comparator( right, right, _preprocessed_right, _preprocessed_right, _has_nulls, _nulls_equal, build); - hash_csr_inclusive_sum(_impl->cumulative_ends.data(), _impl->capacity, stream); + auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, + cudf::get_current_device_resource_ref()}; + CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum( + _impl->cumulative_ends.data(), _impl->cumulative_ends.data(), _impl->capacity, env)); launch_hash_csr_build_fill(right.num_rows(), build_positions.data(), _impl->cumulative_ends.data(), From b057499d32ba5e0fc3095956ceeb7519afc5d3c5 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 19:41:19 +0000 Subject: [PATCH 10/12] Reuse existing utilities in HashCSR join --- cpp/src/join/hash_join/dispatch.cuh | 44 ++++--- cpp/src/join/hash_join/hash_csr.cuh | 76 +++-------- cpp/src/join/hash_join/hash_csr_kernels.cuh | 123 +++++------------- cpp/src/join/hash_join/hash_join.cu | 26 ++-- cpp/src/join/hash_join/hash_join_impl.cuh | 4 +- cpp/src/join/hash_join/match_context.cu | 2 +- .../hash_join/partitioned_join_retrieve.cu | 60 +++------ cpp/src/join/hash_join/retrieve_impl.cuh | 7 +- cpp/src/join/hash_join/size_impl.cuh | 14 +- cpp/src/join/join.cu | 118 +++++++---------- 10 files changed, 173 insertions(+), 301 deletions(-) diff --git a/cpp/src/join/hash_join/dispatch.cuh b/cpp/src/join/hash_join/dispatch.cuh index abe3d5b829e5..f68893c45cf2 100644 --- a/cpp/src/join/hash_join/dispatch.cuh +++ b/cpp/src/join/hash_join/dispatch.cuh @@ -10,43 +10,51 @@ #include #include +#include + #include #include namespace cudf::detail { +/** + * @brief Equality comparator for cuco hash table probing with row-level equality. + */ template -class hash_csr_equal { +class pair_equal { public: - explicit hash_csr_equal(Equal check_row_equality) - : _check_row_equality{std::move(check_row_equality)} - { - } + pair_equal(Equal check_row_equality) : _check_row_equality{std::move(check_row_equality)} {} - __device__ bool operator()(hash_csr_key_type const& lhs, - hash_csr_key_type const& rhs) const noexcept + __device__ __forceinline__ bool operator()( + cuco::pair const& lhs, + cuco::pair const& rhs) const noexcept { using detail::row::lhs_index_type; using detail::row::rhs_index_type; - return _check_row_equality(lhs_index_type{lhs.row}, rhs_index_type{rhs.row}); + + return lhs.first == rhs.first and + _check_row_equality(lhs_index_type{lhs.second}, rhs_index_type{rhs.second}); } private: Equal _check_row_equality; }; -class primitive_hash_csr_equal { +/** + * @brief Equality comparator for cuco hash table probing with primitive row equality. + */ +class primitive_pair_equal { public: - explicit primitive_hash_csr_equal( - cudf::detail::row::primitive::row_equality_comparator check_row_equality) + primitive_pair_equal(cudf::detail::row::primitive::row_equality_comparator check_row_equality) : _check_row_equality{std::move(check_row_equality)} { } - __device__ bool operator()(hash_csr_key_type const& lhs, - hash_csr_key_type const& rhs) const noexcept + __device__ __forceinline__ bool operator()( + cuco::pair const& lhs, + cuco::pair const& rhs) const noexcept { - return _check_row_equality(lhs.row, rhs.row); + return lhs.first == rhs.first and _check_row_equality(lhs.second, rhs.second); } private: @@ -54,7 +62,7 @@ class primitive_hash_csr_equal { }; template -decltype(auto) dispatch_hash_csr_comparator( +decltype(auto) dispatch_join_comparator( table_view const& right_table, table_view const& left_table, std::shared_ptr const& preprocessed_right, @@ -69,7 +77,7 @@ decltype(auto) dispatch_hash_csr_comparator( auto const d_hasher = cudf::detail::row::primitive::row_hasher{left_nulls, preprocessed_left}; auto const d_equal = cudf::detail::row::primitive::row_equality_comparator{ left_nulls, preprocessed_left, preprocessed_right, compare_nulls}; - return std::forward(fn)(primitive_hash_csr_equal{d_equal}, d_hasher); + return std::forward(fn)(primitive_pair_equal{d_equal}, d_hasher); } auto const d_hasher = @@ -79,11 +87,11 @@ decltype(auto) dispatch_hash_csr_comparator( if (cudf::detail::has_nested_columns(left_table)) { auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); - return std::forward(fn)(hash_csr_equal{d_equal}, d_hasher); + return std::forward(fn)(pair_equal{d_equal}, d_hasher); } auto const d_equal = row_comparator.equal_to(left_nulls, compare_nulls); - return std::forward(fn)(hash_csr_equal{d_equal}, d_hasher); + return std::forward(fn)(pair_equal{d_equal}, d_hasher); } } // namespace cudf::detail diff --git a/cpp/src/join/hash_join/hash_csr.cuh b/cpp/src/join/hash_join/hash_csr.cuh index 4be73a5efe64..f44718eaa347 100644 --- a/cpp/src/join/hash_join/hash_csr.cuh +++ b/cpp/src/join/hash_join/hash_csr.cuh @@ -5,75 +5,35 @@ #pragma once +#include #include #include +#include #include #include -#include namespace cudf::detail { -struct hash_csr_key_type { - hash_value_type hash; - size_type row; -}; - -constexpr size_type hash_csr_empty_slot = size_type{-1}; -constexpr std::uint64_t hash_csr_empty_entry = std::numeric_limits::max(); -constexpr std::uint64_t hash_csr_empty_build_position = std::numeric_limits::max(); - -static __device__ std::uint64_t pack_hash_csr_entry(hash_value_type hash, size_type row) -{ - return (static_cast(hash) << 32) | static_cast(row); -} - -static __device__ hash_value_type unpack_hash_csr_hash(std::uint64_t value) -{ - return static_cast(value >> 32); -} - -static __device__ size_type unpack_hash_csr_row(std::uint64_t value) -{ - return static_cast(static_cast(value)); -} - -static __device__ std::uint64_t pack_hash_csr_build_position(std::uint32_t slot, size_type rank) -{ - return (static_cast(slot) << 32) | static_cast(rank); -} - -static __device__ std::uint32_t unpack_hash_csr_build_slot(std::uint64_t value) -{ - return static_cast(value >> 32); -} - -static __device__ size_type unpack_hash_csr_build_rank(std::uint64_t value) -{ - return static_cast(static_cast(value)); -} +using hash_csr_key_type = cuco::pair; +using hash_csr_build_position_type = cuco::pair; struct hash_csr_map_view { - std::uint64_t* entries; + hash_csr_key_type* entries; std::uint32_t capacity; std::uint32_t mask; template __device__ std::uint32_t insert(hash_csr_key_type key, Equal equal) const { - auto const desired = pack_hash_csr_entry(key.hash, key.row); for (std::uint32_t step = 0; step < capacity; ++step) { - auto const slot = (static_cast(key.hash) + step) & mask; - auto entry_ref = cuda::atomic_ref{entries[slot]}; - auto old = hash_csr_empty_entry; - if (entry_ref.compare_exchange_strong(old, desired, cuda::memory_order_relaxed)) { - return slot; - } - if (unpack_hash_csr_hash(old) == key.hash && - equal(key, hash_csr_key_type{unpack_hash_csr_hash(old), unpack_hash_csr_row(old)})) { - return slot; - } + auto const slot = (static_cast(key.first) + step) & mask; + auto entry_ref = + cuda::atomic_ref{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; } @@ -82,14 +42,10 @@ struct hash_csr_map_view { __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(key.hash) + step) & mask; + auto const slot = (static_cast(key.first) + step) & mask; auto const current = entries[slot]; - if (current == hash_csr_empty_entry) { return capacity; } - if (unpack_hash_csr_hash(current) == key.hash && - equal(key, - hash_csr_key_type{unpack_hash_csr_hash(current), unpack_hash_csr_row(current)})) { - return slot; - } + if (current.second == CUDF_SIZE_TYPE_SENTINEL) { return capacity; } + if (equal(key, current)) { return slot; } } return capacity; } @@ -104,9 +60,7 @@ struct hash_csr_view { return slot == 0 ? size_type{0} : cumulative_ends[slot - 1]; } - __device__ size_type end(size_type slot) const { return cumulative_ends[slot]; } - - __device__ size_type size(size_type slot) const { return end(slot) - begin(slot); } + __device__ size_type size(size_type slot) const { return cumulative_ends[slot] - begin(slot); } }; } // namespace cudf::detail diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index 3d753e8250e0..fba8c0f1f35b 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -7,27 +7,17 @@ #include "hash_csr.cuh" -#include -#include -#include #include #include #include +#include #include #include -#include -#include -#include -#include -#include -#include -#include -#include +#include +#include -#include #include -#include namespace cudf::detail { @@ -39,7 +29,7 @@ constexpr thread_index_type hash_csr_outputs_per_lane = 32; template CUDF_KERNEL void hash_csr_build_count_kernel(size_type num_rows, bitmask_type const* valid_rows, - std::uint64_t* build_positions, + hash_csr_build_position_type* build_positions, size_type* slot_counts, hash_csr_map_view map, Equal equal, @@ -49,23 +39,23 @@ CUDF_KERNEL void hash_csr_build_count_kernel(size_type num_rows, for (auto row = grid_1d::global_thread_id(); row < num_rows; row += stride) { auto const index = static_cast(row); if (valid_rows != nullptr && !cudf::bit_is_set(valid_rows, index)) { - build_positions[index] = hash_csr_empty_build_position; + build_positions[index] = {std::uint32_t{-1}, CUDF_SIZE_TYPE_SENTINEL}; continue; } auto const slot = map.insert(hash_csr_key_type{hasher(index), index}, equal); if (slot == map.capacity) { - build_positions[index] = hash_csr_empty_build_position; + build_positions[index] = {std::uint32_t{-1}, CUDF_SIZE_TYPE_SENTINEL}; continue; } auto slot_count_ref = cuda::atomic_ref{slot_counts[slot]}; auto const rank = slot_count_ref.fetch_add(size_type{1}, cuda::memory_order_relaxed); - build_positions[index] = pack_hash_csr_build_position(slot, rank); + build_positions[index] = {slot, rank}; } } CUDF_KERNEL void hash_csr_build_fill_kernel(size_type num_rows, - std::uint64_t const* build_positions, + hash_csr_build_position_type const* build_positions, size_type const* cumulative_ends, size_type* values) { @@ -73,9 +63,9 @@ CUDF_KERNEL void hash_csr_build_fill_kernel(size_type num_rows, for (auto row = grid_1d::global_thread_id(); row < num_rows; row += stride) { auto const index = static_cast(row); auto const position = build_positions[index]; - if (position == hash_csr_empty_build_position) { continue; } - auto const slot = unpack_hash_csr_build_slot(position); - auto const rank = unpack_hash_csr_build_rank(position); + if (position.first == std::uint32_t{-1}) { continue; } + auto const slot = position.first; + auto const rank = position.second; auto const begin = slot == 0 ? size_type{0} : cumulative_ends[slot - 1]; values[begin + rank] = index; } @@ -104,7 +94,7 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, auto const found = slot != map.capacity; auto const count = found ? csr.size(static_cast(slot)) : size_type{0}; if (probe_slots != nullptr) { - probe_slots[index] = found ? static_cast(slot) : hash_csr_empty_slot; + probe_slots[index] = found ? static_cast(slot) : CUDF_SIZE_TYPE_SENTINEL; } if (match_counts != nullptr) { match_counts[index] = is_outer ? cuda::std::max(count, size_type{1}) : count; @@ -126,7 +116,7 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, template void launch_hash_csr_build_count(size_type num_rows, bitmask_type const* valid_rows, - std::uint64_t* build_positions, + hash_csr_build_position_type* build_positions, size_type* slot_counts, hash_csr_map_view map, Equal equal, @@ -143,11 +133,12 @@ void launch_hash_csr_build_count(size_type num_rows, CUDF_CUDA_TRY(cudaGetLastError()); } -[[maybe_unused]] static void launch_hash_csr_build_fill(size_type num_rows, - std::uint64_t const* build_positions, - size_type const* cumulative_ends, - size_type* values, - rmm::cuda_stream_view stream) +[[maybe_unused]] static void launch_hash_csr_build_fill( + size_type num_rows, + hash_csr_build_position_type const* build_positions, + size_type const* cumulative_ends, + size_type* values, + rmm::cuda_stream_view stream) { if (num_rows == 0) { return; } auto const config = grid_1d{num_rows, hash_csr_block_size}; @@ -188,59 +179,6 @@ void launch_hash_csr_probe_count(size_type num_rows, CUDF_CUDA_TRY(cudaGetLastError()); } -[[maybe_unused]] static std::int64_t hash_csr_scan_counts(size_type const* counts, - size_type num_rows, - std::int64_t* offsets, - rmm::cuda_stream_view stream) -{ - auto const mr = cudf::get_current_device_resource_ref(); - cudf::detail::device_scalar output_size(stream, mr); - auto output = cudf::detail::make_sizes_to_offsets_iterator( - offsets, offsets + num_rows + 1, output_size.data()); - auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, mr}; - CUDF_CUDA_TRY(cub::DeviceScan::ExclusiveScan( - counts, output, cuda::std::plus<>{}, std::int64_t{0}, num_rows + 1, env)); - return output_size.value(stream); -} - -struct hash_csr_count_to_int64 { - __device__ std::int64_t operator()(size_type value) const - { - return static_cast(value); - } -}; - -[[maybe_unused]] static std::int64_t hash_csr_reduce_counts(size_type const* counts, - size_type num_rows, - rmm::cuda_stream_view stream) -{ - if (num_rows == 0) { return 0; } - auto const mr = cudf::get_current_device_resource_ref(); - auto input = cuda::transform_iterator{counts, hash_csr_count_to_int64{}}; - cudf::detail::device_scalar result(stream, mr); - auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, mr}; - CUDF_CUDA_TRY(cub::DeviceReduce::Sum(input, result.data(), num_rows, env)); - return result.value(stream); -} - -static __device__ size_type hash_csr_find_probe_row_in_range(std::int64_t const* offsets, - size_type first_probe, - size_type last_probe, - std::int64_t output_index) -{ - auto first = static_cast(first_probe); - auto last = static_cast(last_probe) + 2; - while (first < last) { - auto const middle = first + (last - first) / 2; - if (offsets[middle] <= output_index) { - first = middle + 1; - } else { - last = middle; - } - } - return static_cast(first - 1); -} - template CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, size_type num_probe_rows, @@ -252,7 +190,9 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, size_type* left_indices, size_type* right_indices) { - auto const lane_id = static_cast(threadIdx.x) % cudf::detail::warp_size; + auto const warp = cooperative_groups::tiled_partition( + cooperative_groups::this_thread_block()); + auto const lane_id = static_cast(warp.thread_rank()); auto const warp_in_block = static_cast(threadIdx.x) / cudf::detail::warp_size; auto const global_warp = static_cast(blockIdx.x) * hash_csr_warps_per_block + warp_in_block; @@ -263,10 +203,11 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, size_type endpoint_probe{}; if (lane_id < 2) { auto const endpoint = lane_id == 0 ? range_begin : range_end - 1; - endpoint_probe = hash_csr_find_probe_row_in_range(offsets, 0, num_probe_rows - 1, endpoint); + endpoint_probe = static_cast( + cuda::std::upper_bound(offsets, offsets + num_probe_rows + 1, endpoint) - offsets - 1); } - auto const first_probe = __shfl_sync(0xffffffff, endpoint_probe, 0); - auto const last_probe = __shfl_sync(0xffffffff, endpoint_probe, 1); + auto const first_probe = warp.shfl(endpoint_probe, 0); + auto const last_probe = warp.shfl(endpoint_probe, 1); #pragma unroll for (thread_index_type item = 0; item < hash_csr_outputs_per_lane; ++item) { @@ -275,11 +216,14 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, auto const probe_row = first_probe == last_probe ? first_probe - : hash_csr_find_probe_row_in_range(offsets, first_probe, last_probe, output_index); + : static_cast(cuda::std::upper_bound(offsets + first_probe, + offsets + last_probe + 2, + output_index) - + offsets - 1); auto const slot = probe_slots[probe_row]; left_indices[output_index] = probe_row + left_index_offset; if constexpr (is_outer) { - if (slot == hash_csr_empty_slot) { + if (slot == CUDF_SIZE_TYPE_SENTINEL) { right_indices[output_index] = JoinNoMatch; continue; } @@ -305,11 +249,12 @@ void launch_hash_csr_retrieve(std::int64_t output_size, auto const min_blocks = size_type{2} * cudf::detail::num_multiprocessors(); constexpr auto outputs_per_block = hash_csr_warps_per_block * cudf::detail::warp_size * hash_csr_outputs_per_lane; - auto const requested_blocks = (output_size + outputs_per_block - 1) / outputs_per_block; + auto const requested_blocks = + cudf::util::div_rounding_up_safe(output_size, static_cast(outputs_per_block)); auto const num_blocks = static_cast(cuda::std::max(requested_blocks, min_blocks)); auto const num_warps = static_cast(num_blocks) * hash_csr_warps_per_block; - auto const outputs_per_warp = (output_size + num_warps - 1) / num_warps; + auto const outputs_per_warp = cudf::util::div_rounding_up_safe(output_size, num_warps); hash_csr_retrieve_kernel <<>>(output_size, diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index 905a38578503..e51b7932f486 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -9,7 +9,6 @@ #include "join/join_common_utils.cuh" #include -#include #include #include #include @@ -21,7 +20,7 @@ #include #include -#include +#include #include #include @@ -64,13 +63,13 @@ void validate_hash_join_probe(table_view const& right, table_view const& left, b namespace { std::uint32_t hash_csr_capacity(size_type rows, double load_factor) { - auto const checked = checked_load_factor(load_factor); - auto const requested = std::max(static_cast(rows) + 1, + auto const checked = checked_load_factor(load_factor); + auto const requested = std::max(static_cast(rows) + 1, std::ceil(static_cast(rows) / checked)); - std::uint64_t capacity = 1; - while (static_cast(capacity) < requested) { - capacity <<= 1; - } + CUDF_EXPECTS(requested <= std::numeric_limits::max(), + "HashCSR table capacity is not representable", + std::overflow_error); + auto const capacity = cuda::std::bit_ceil(static_cast(requested)); CUDF_EXPECTS(capacity <= std::numeric_limits::max(), "HashCSR table capacity is not representable", std::overflow_error); @@ -107,8 +106,10 @@ hash_join::hash_join(cudf::table_view const& right, CUDF_EXPECTS(0 != right.num_columns(), "Hash join right table is empty", std::invalid_argument); if (_is_empty) { return; } - CUDF_CUDA_TRY(cudaMemsetAsync( - _impl->entries.data(), 0xff, _impl->entries.size() * sizeof(std::uint64_t), stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(_impl->entries.data(), + 0xff, + _impl->entries.size() * sizeof(hash_csr_key_type), + stream.value())); CUDF_CUDA_TRY(cudaMemsetAsync(_impl->cumulative_ends.data(), 0, _impl->cumulative_ends.size() * sizeof(size_type), @@ -119,7 +120,8 @@ hash_join::hash_join(cudf::table_view const& right, auto const valid_rows = _nulls_equal == null_equality::UNEQUAL ? static_cast(row_bitmask.data()) : nullptr; - rmm::device_uvector build_positions(right.num_rows(), stream, temp_mr); + rmm::device_uvector build_positions( + right.num_rows(), stream, temp_mr); auto build = [&](auto equality, auto hasher) { launch_hash_csr_build_count(right.num_rows(), valid_rows, @@ -130,7 +132,7 @@ hash_join::hash_join(cudf::table_view const& right, hasher, stream); }; - dispatch_hash_csr_comparator( + dispatch_join_comparator( right, right, _preprocessed_right, _preprocessed_right, _has_nulls, _nulls_equal, build); auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, cudf::get_current_device_resource_ref()}; diff --git a/cpp/src/join/hash_join/hash_join_impl.cuh b/cpp/src/join/hash_join/hash_join_impl.cuh index 96d0152eab34..fbcf3953f704 100644 --- a/cpp/src/join/hash_join/hash_join_impl.cuh +++ b/cpp/src/join/hash_join/hash_join_impl.cuh @@ -32,13 +32,13 @@ struct hash_join::impl { hash_csr_map_view map_view() const { - return {const_cast(entries.data()), capacity, capacity - 1}; + return {const_cast(entries.data()), capacity, capacity - 1}; } hash_csr_view csr_view() const { return {cumulative_ends.data(), values.data()}; } cuda::mr::any_resource _mr; - rmm::device_uvector entries; + rmm::device_uvector entries; rmm::device_uvector cumulative_ends; rmm::device_uvector values; std::uint32_t capacity; diff --git a/cpp/src/join/hash_join/match_context.cu b/cpp/src/join/hash_join/match_context.cu index 11b42848a381..9e41bda9c728 100644 --- a/cpp/src/join/hash_join/match_context.cu +++ b/cpp/src/join/hash_join/match_context.cu @@ -74,7 +74,7 @@ std::unique_ptr> hash_join::make_match_co } }; - dispatch_hash_csr_comparator( + dispatch_join_comparator( _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); return match_counts; diff --git a/cpp/src/join/hash_join/partitioned_join_retrieve.cu b/cpp/src/join/hash_join/partitioned_join_retrieve.cu index 9c0d3b365624..eb081957c464 100644 --- a/cpp/src/join/hash_join/partitioned_join_retrieve.cu +++ b/cpp/src/join/hash_join/partitioned_join_retrieve.cu @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -21,39 +22,10 @@ #include #include -#include -#include -#include +#include +#include namespace cudf::detail { -namespace { - -/** - * @brief Returns trivial left/right index pairs for an outer join when the build side is empty. - */ -std::pair>, - std::unique_ptr>> -make_trivial_outer_indices(size_type left_start_idx, - size_type partition_size, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) -{ - auto left_indices = std::make_unique>(partition_size, stream, mr); - auto right_indices = std::make_unique>(partition_size, stream, mr); - auto out = cuda::zip_iterator(left_indices->begin(), right_indices->begin()); - thrust::tabulate(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - out, - out + partition_size, - cuda::proclaim_return_type>( - [left_start_idx] __device__(auto i) { - return cuda::std::tuple{static_cast(left_start_idx + i), - JoinNoMatch}; - })); - return std::pair(std::move(left_indices), std::move(right_indices)); -} - -} // namespace - template std::pair>, std::unique_ptr>> @@ -97,9 +69,15 @@ hash_join::partitioned_join_retrieve(join_kind join, if (join == join_kind::INNER_JOIN) { return std::pair(std::make_unique>(0, stream, mr), std::make_unique>(0, stream, mr)); - } else { - return make_trivial_outer_indices(left_start_idx, partition_size, stream, mr); } + auto left_indices = + std::make_unique>(partition_size, stream, mr); + auto right_indices = + std::make_unique>(partition_size, stream, mr); + auto const exec = rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()); + thrust::sequence(exec, left_indices->begin(), left_indices->end(), left_start_idx); + thrust::fill(exec, right_indices->begin(), right_indices->end(), JoinNoMatch); + return std::pair(std::move(left_indices), std::move(right_indices)); } // Slice the left table to the partition range @@ -122,7 +100,7 @@ hash_join::partitioned_join_retrieve(join_kind join, auto offsets = cudf::detail::make_zeroed_device_uvector_async( static_cast(partition_size) + 1, stream, temp_mr); auto const output_size = - hash_csr_scan_counts(counts.data(), partition_size, offsets.data(), stream); + cudf::detail::sizes_to_offsets(counts.begin(), counts.end(), offsets.begin(), 0, stream); CUDF_EXPECTS(output_size >= 0, "Join output size overflowed", std::overflow_error); rmm::device_uvector probe_slots(partition_size, stream, temp_mr); @@ -157,13 +135,13 @@ hash_join::partitioned_join_retrieve(join_kind join, stream); } }; - dispatch_hash_csr_comparator(_right, - left_partition_view, - _preprocessed_right, - preprocessed_left, - _has_nulls, - _nulls_equal, - save_slots); + dispatch_join_comparator(_right, + left_partition_view, + _preprocessed_right, + preprocessed_left, + _has_nulls, + _nulls_equal, + save_slots); auto left_indices = std::make_unique>( static_cast(output_size), stream, mr); diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index d98507546216..95a2a8c63299 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -11,6 +11,7 @@ #include #include +#include #include #include #include @@ -87,13 +88,13 @@ hash_join::join_retrieve(cudf::table_view const& left, stream); } }; - dispatch_hash_csr_comparator( + dispatch_join_comparator( _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); auto offsets = cudf::detail::make_zeroed_device_uvector_async( static_cast(left.num_rows()) + 1, stream, temp_mr); - auto const actual_size = - hash_csr_scan_counts(match_counts.data(), left.num_rows(), offsets.data(), stream); + auto const actual_size = cudf::detail::sizes_to_offsets( + match_counts.begin(), match_counts.end(), offsets.begin(), 0, stream); CUDF_EXPECTS(actual_size >= 0, "Join output size overflowed", std::overflow_error); auto const join_size = Join != join_kind::FULL_JOIN && output_size.has_value() ? *output_size diff --git a/cpp/src/join/hash_join/size_impl.cuh b/cpp/src/join/hash_join/size_impl.cuh index 972ba8ba0a9f..a0cc1a8e0920 100644 --- a/cpp/src/join/hash_join/size_impl.cuh +++ b/cpp/src/join/hash_join/size_impl.cuh @@ -8,10 +8,13 @@ #include "dispatch.cuh" #include "hash_csr_kernels.cuh" +#include #include #include #include +#include + namespace cudf::detail { template @@ -57,9 +60,10 @@ std::size_t hash_join::join_size(cudf::table_view const& left, hasher, stream); }; - dispatch_hash_csr_comparator( + dispatch_join_comparator( _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); - auto const output_size = hash_csr_reduce_counts(match_counts.data(), left.num_rows(), stream); + auto const output_size = cudf::detail::reduce( + match_counts.begin(), match_counts.end(), std::int64_t{0}, cuda::std::plus<>{}, stream); CUDF_EXPECTS(output_size >= 0, "Join output size overflowed", std::overflow_error); return static_cast(output_size); } @@ -104,11 +108,11 @@ std::size_t hash_join::join_size(cudf::table_view const& left, hasher, stream); }; - dispatch_hash_csr_comparator( + dispatch_join_comparator( _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches); - auto const left_output_size = - hash_csr_reduce_counts(match_counts.data(), left.num_rows(), stream); + auto const left_output_size = cudf::detail::reduce( + match_counts.begin(), match_counts.end(), std::int64_t{0}, cuda::std::plus<>{}, stream); auto const matched_right_rows = matched_build_rows.value(stream); CUDF_EXPECTS(left_output_size >= 0, "Join output size overflowed", std::overflow_error); auto const output_size = static_cast(left_output_size) + diff --git a/cpp/src/join/join.cu b/cpp/src/join/join.cu index f8f788d2b848..9334125e05df 100644 --- a/cpp/src/join/join.cu +++ b/cpp/src/join/join.cu @@ -1,12 +1,11 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ #include "join_common_utils.hpp" #include #include -#include #include #include #include @@ -20,21 +19,11 @@ #include #include -#include #include namespace cudf { namespace detail { -namespace { -bool has_dictionary_columns(table_view const& table) -{ - return std::any_of(table.begin(), table.end(), [](column_view const& column) { - return column.type().id() == type_id::DICTIONARY32; - }); -} -} // namespace - std::pair>, std::unique_ptr>> inner_join(table_view const& left_input, @@ -43,34 +32,31 @@ inner_join(table_view const& left_input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto join_tables = [&](table_view const& left, table_view const& right) { - auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) - ? cudf::nullable_join::YES - : cudf::nullable_join::NO; - if (right.num_rows() > left.num_rows()) { - cudf::hash_join::impl_type hj_obj(left, - has_nulls == cudf::nullable_join::YES, - compare_nulls, - CUCO_DESIRED_LOAD_FACTOR, - stream, - cudf::get_current_device_resource_ref()); - auto [right_result, left_result] = hj_obj.inner_join(right, std::nullopt, stream, mr); - return std::pair(std::move(left_result), std::move(right_result)); - } - cudf::hash_join::impl_type hj_obj(right, - has_nulls == cudf::nullable_join::YES, - compare_nulls, - CUCO_DESIRED_LOAD_FACTOR, - stream, - cudf::get_current_device_resource_ref()); - return hj_obj.inner_join(left, std::nullopt, stream, mr); - }; - - if (!has_dictionary_columns(left_input)) { return join_tables(left_input, right_input); } - + // Make sure any dictionary columns have matched key sets. + // This will return any new dictionary columns created as well as updated table_views. auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, stream, cudf::get_current_device_resource_ref()); - return join_tables(matched.second.front(), matched.second.back()); + {left_input, right_input}, + stream, + cudf::get_current_device_resource_ref()); // temporary objects returned + + // now rebuild the table views with the updated ones + auto const left = matched.second.front(); + auto const right = matched.second.back(); + auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) + ? cudf::nullable_join::YES + : cudf::nullable_join::NO; + + // For `inner_join`, we can freely choose either the `left` or `right` table to use for + // building/probing the hash map. Because building is typically more expensive than probing, we + // build the hash map from the smaller table. + if (right.num_rows() > left.num_rows()) { + cudf::hash_join hj_obj(left, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); + auto [right_result, left_result] = hj_obj.inner_join(right, std::nullopt, stream, mr); + return std::pair(std::move(left_result), std::move(right_result)); + } else { + cudf::hash_join hj_obj(right, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); + return hj_obj.inner_join(left, std::nullopt, stream, mr); + } } std::pair>, @@ -81,24 +67,21 @@ left_join(table_view const& left_input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto join_tables = [&](table_view const& left, table_view const& right) { - auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) + // Make sure any dictionary columns have matched key sets. + // This will return any new dictionary columns created as well as updated table_views. + auto matched = cudf::dictionary::detail::match_dictionaries( + {left_input, right_input}, // these should match + stream, + cudf::get_current_device_resource_ref()); // temporary objects returned + // now rebuild the table views with the updated ones + table_view const left = matched.second.front(); + table_view const right = matched.second.back(); + auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) ? cudf::nullable_join::YES : cudf::nullable_join::NO; - cudf::hash_join::impl_type hj_obj(right, - has_nulls == cudf::nullable_join::YES, - compare_nulls, - CUCO_DESIRED_LOAD_FACTOR, - stream, - cudf::get_current_device_resource_ref()); - return hj_obj.left_join(left, std::nullopt, stream, mr); - }; - if (!has_dictionary_columns(left_input)) { return join_tables(left_input, right_input); } - - auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, stream, cudf::get_current_device_resource_ref()); - return join_tables(matched.second.front(), matched.second.back()); + cudf::hash_join hj_obj(right, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); + return hj_obj.left_join(left, std::nullopt, stream, mr); } std::pair>, @@ -109,24 +92,21 @@ full_join(table_view const& left_input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - auto join_tables = [&](table_view const& left, table_view const& right) { - auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) + // Make sure any dictionary columns have matched key sets. + // This will return any new dictionary columns created as well as updated table_views. + auto matched = cudf::dictionary::detail::match_dictionaries( + {left_input, right_input}, // these should match + stream, + cudf::get_current_device_resource_ref()); // temporary objects returned + // now rebuild the table views with the updated ones + table_view const left = matched.second.front(); + table_view const right = matched.second.back(); + auto const has_nulls = cudf::has_nested_nulls(left) || cudf::has_nested_nulls(right) ? cudf::nullable_join::YES : cudf::nullable_join::NO; - cudf::hash_join::impl_type hj_obj(right, - has_nulls == cudf::nullable_join::YES, - compare_nulls, - CUCO_DESIRED_LOAD_FACTOR, - stream, - cudf::get_current_device_resource_ref()); - return hj_obj.full_join(left, std::nullopt, stream, mr); - }; - - if (!has_dictionary_columns(left_input)) { return join_tables(left_input, right_input); } - auto matched = cudf::dictionary::detail::match_dictionaries( - {left_input, right_input}, stream, cudf::get_current_device_resource_ref()); - return join_tables(matched.second.front(), matched.second.back()); + cudf::hash_join hj_obj(right, has_nulls, compare_nulls, CUCO_DESIRED_LOAD_FACTOR, stream); + return hj_obj.full_join(left, std::nullopt, stream, mr); } } // namespace detail From a9f4993569e84bfd5d9d617507f766697e78fb14 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 20:55:05 +0000 Subject: [PATCH 11/12] Use explicit CUB storage for HashCSR scan --- cpp/src/join/hash_join/hash_join.cu | 20 ++++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/cpp/src/join/hash_join/hash_join.cu b/cpp/src/join/hash_join/hash_join.cu index e51b7932f486..828ae6e31b11 100644 --- a/cpp/src/join/hash_join/hash_join.cu +++ b/cpp/src/join/hash_join/hash_join.cu @@ -20,6 +20,8 @@ #include #include +#include + #include #include @@ -134,10 +136,20 @@ hash_join::hash_join(cudf::table_view const& right, }; dispatch_join_comparator( right, right, _preprocessed_right, _preprocessed_right, _has_nulls, _nulls_equal, build); - auto env = cuda::std::execution::env{cuda::stream_ref{stream.value()}, - cudf::get_current_device_resource_ref()}; - CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum( - _impl->cumulative_ends.data(), _impl->cumulative_ends.data(), _impl->capacity, env)); + std::size_t temp_storage_bytes{}; + CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum(nullptr, + temp_storage_bytes, + _impl->cumulative_ends.data(), + _impl->cumulative_ends.data(), + _impl->capacity, + stream.value())); + rmm::device_buffer temp_storage(temp_storage_bytes, stream, temp_mr); + CUDF_CUDA_TRY(cub::DeviceScan::InclusiveSum(temp_storage.data(), + temp_storage_bytes, + _impl->cumulative_ends.data(), + _impl->cumulative_ends.data(), + _impl->capacity, + stream.value())); launch_hash_csr_build_fill(right.num_rows(), build_positions.data(), _impl->cumulative_ends.data(), From dcef5dd2aec143cf65248a6ea695b252c88d1ca2 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Wed, 12 Aug 2026 21:48:56 +0000 Subject: [PATCH 12/12] Clean up HashCSR join dispatch --- cpp/include/cudf/detail/join/hash_join.hpp | 2 +- cpp/src/join/hash_join/hash_csr_kernels.cuh | 16 ++++----- cpp/src/join/hash_join/retrieve_impl.cuh | 36 +++++++-------------- 3 files changed, 20 insertions(+), 34 deletions(-) diff --git a/cpp/include/cudf/detail/join/hash_join.hpp b/cpp/include/cudf/detail/join/hash_join.hpp index 4a04234b8d70..0d311d2a69bc 100644 --- a/cpp/include/cudf/detail/join/hash_join.hpp +++ b/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 the build-side 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 diff --git a/cpp/src/join/hash_join/hash_csr_kernels.cuh b/cpp/src/join/hash_join/hash_csr_kernels.cuh index fba8c0f1f35b..11d9398303d5 100644 --- a/cpp/src/join/hash_join/hash_csr_kernels.cuh +++ b/cpp/src/join/hash_join/hash_csr_kernels.cuh @@ -71,7 +71,7 @@ CUDF_KERNEL void hash_csr_build_fill_kernel(size_type num_rows, } } -template +template CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, bitmask_type const* valid_rows, size_type* probe_slots, @@ -97,7 +97,7 @@ CUDF_KERNEL void hash_csr_probe_count_kernel(size_type num_rows, probe_slots[index] = found ? static_cast(slot) : CUDF_SIZE_TYPE_SENTINEL; } if (match_counts != nullptr) { - match_counts[index] = is_outer ? cuda::std::max(count, size_type{1}) : count; + match_counts[index] = IsOuter ? cuda::std::max(count, size_type{1}) : count; } if (found && matched_slots != nullptr) { @@ -150,7 +150,7 @@ void launch_hash_csr_build_count(size_type num_rows, CUDF_CUDA_TRY(cudaGetLastError()); } -template +template void launch_hash_csr_probe_count(size_type num_rows, bitmask_type const* valid_rows, size_type* probe_slots, @@ -165,7 +165,7 @@ void launch_hash_csr_probe_count(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 + hash_csr_probe_count_kernel <<>>(num_rows, valid_rows, probe_slots, @@ -179,7 +179,7 @@ void launch_hash_csr_probe_count(size_type num_rows, CUDF_CUDA_TRY(cudaGetLastError()); } -template +template CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, size_type num_probe_rows, std::int64_t outputs_per_warp, @@ -222,7 +222,7 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, offsets - 1); auto const slot = probe_slots[probe_row]; left_indices[output_index] = probe_row + left_index_offset; - if constexpr (is_outer) { + if constexpr (IsOuter) { if (slot == CUDF_SIZE_TYPE_SENTINEL) { right_indices[output_index] = JoinNoMatch; continue; @@ -234,7 +234,7 @@ CUDF_KERNEL void hash_csr_retrieve_kernel(std::int64_t output_size, } } -template +template void launch_hash_csr_retrieve(std::int64_t output_size, size_type num_probe_rows, std::int64_t const* offsets, @@ -256,7 +256,7 @@ void launch_hash_csr_retrieve(std::int64_t output_size, auto const num_warps = static_cast(num_blocks) * hash_csr_warps_per_block; auto const outputs_per_warp = cudf::util::div_rounding_up_safe(output_size, num_warps); - hash_csr_retrieve_kernel + hash_csr_retrieve_kernel <<>>(output_size, num_probe_rows, outputs_per_warp, diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index 95a2a8c63299..0efe6998f581 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -62,31 +62,17 @@ hash_join::join_retrieve(cudf::table_view const& left, : nullptr; auto count_matches = [&](auto equality, auto hasher) { - if constexpr (Join == join_kind::INNER_JOIN) { - launch_hash_csr_probe_count(left.num_rows(), - valid_rows, - probe_slots.data(), - match_counts.data(), - nullptr, - nullptr, - _impl->map_view(), - _impl->csr_view(), - equality, - hasher, - stream); - } else { - launch_hash_csr_probe_count(left.num_rows(), - valid_rows, - probe_slots.data(), - match_counts.data(), - nullptr, - nullptr, - _impl->map_view(), - _impl->csr_view(), - equality, - hasher, - stream); - } + launch_hash_csr_probe_count(left.num_rows(), + valid_rows, + probe_slots.data(), + match_counts.data(), + nullptr, + nullptr, + _impl->map_view(), + _impl->csr_view(), + equality, + hasher, + stream); }; dispatch_join_comparator( _right, left, _preprocessed_right, preprocessed_left, _has_nulls, _nulls_equal, count_matches);