Add specialized dispatch to improve occupancy for hash table operations in distinct join #16321
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
bdice
reviewed
Jul 23, 2024
| probing_scheme_type, | ||
| cudf::detail::cuco_allocator, | ||
| cuco_storage_type>; | ||
| using hash_table_type = std::variant< |
There was a problem hiding this comment.
Could we simplify this with something like:
template<typename Comparator>
using static_set_with_comparator = cuco::static_set<cuco::pair<hash_value_type, rhs_index_type>,
cuco::extent<size_type>,
cuda::thread_scope_device,
comparator_adapter<cudf::experimental::row::equality::
strong_index_comparator_adapter<Comparator>>,
probing_scheme_type,
cudf::detail::cuco_allocator,
cuco_storage_type>;
using hash_table_type = std::variant<
static_set_with_comparator<row_comparator>,
static_set_with_comparator<row_comparator_no_nested>,
static_set_with_comparator<row_comparator_no_compound>>;| /** | ||
| * @brief Compares the specified elements for equality. | ||
| * | ||
| * is_equality_comparable differs from implementation for std::equality_comparable and considers | ||
| * void as and equality comparable type. Thus we need to disable this for when type is void. |
There was a problem hiding this comment.
Suggested change
| * void as and equality comparable type. Thus we need to disable this for when type is void. | |
| * void as an equality comparable type. Thus we need to disable this for when type is void. |
cpp/src/join/distinct_hash_join.cu
Outdated
| return std::visit( | ||
| [&](auto& comparator) { | ||
| return ret_type{ | ||
| std::in_place_type< |
There was a problem hiding this comment.
Can this use in_place_type_t and drop the ::type?
cpp/src/join/distinct_hash_join.cu
Outdated
| cuco::static_set<cuco::pair<hash_value_type, rhs_index_type>, | ||
| cuco::extent<size_type>, | ||
| cuda::thread_scope_device, | ||
| typename std::remove_reference<decltype(comparator_adapter)>::type, |
There was a problem hiding this comment.
Suggested change
| typename std::remove_reference<decltype(comparator_adapter)>::type, | |
| typename std::remove_reference_t<decltype(comparator_adapter)>, |
cpp/src/join/distinct_hash_join.cu
Outdated
| [&](auto&& hasher, auto&& hash_table) { | ||
| auto const iter = cudf::detail::make_counting_transform_iterator( | ||
| 0, | ||
| build_keys_fn<typename std::remove_reference<decltype(hasher)>::type, rhs_index_type>{ |
There was a problem hiding this comment.
Suggested change
| build_keys_fn<typename std::remove_reference<decltype(hasher)>::type, rhs_index_type>{ | |
| build_keys_fn<typename std::remove_reference_t<decltype(hasher)>, rhs_index_type>{ |
cpp/src/join/distinct_hash_join.cu
Outdated
| [&](auto&& hasher, auto&& hash_table) { | ||
| auto const iter = cudf::detail::make_counting_transform_iterator( | ||
| 0, | ||
| build_keys_fn<typename std::remove_reference<decltype(hasher)>::type, lhs_index_type>{ |
There was a problem hiding this comment.
Suggested change
| build_keys_fn<typename std::remove_reference<decltype(hasher)>::type, lhs_index_type>{ | |
| build_keys_fn<typename std::remove_reference_t<decltype(hasher)>, lhs_index_type>{ |
| // used to circumvent conflicts between arrays of different types between | ||
| // different template instantiations due to the extern specifier. | ||
| extern __shared__ char raw_intermediate_storage[]; | ||
| cudf::ast::detail::IntermediateDataType<has_nulls>* intermediate_storage = |
There was a problem hiding this comment.
Can this be auto since you're reinterpret-casting it to the appropriate type?
| cudf::size_type const right_num_rows = right_table.num_rows(); | ||
| auto const outer_num_rows = left_num_rows; | ||
|
|
||
| cudf::size_type outer_row_index = threadIdx.x + blockIdx.x * block_size; |
There was a problem hiding this comment.
We avoid raw threadIdx / blockIdx math. We have utility functions for abstracting this logic and ensuring safe types are used:
See examples:
There was a problem hiding this comment.
I only separated out the instantiation of the kernels in the MR, and copied the previous implementation. Unsure if I should address it in this MR, maybe a followup would be better?
| #include <cudf/utilities/span.hpp> | ||
|
|
||
| #include <cub/cub.cuh> | ||
| #include "mixed_join_kernel_semi_impl.cuh" |
There was a problem hiding this comment.
We should be consistent on including the directory name. This file is next to mixed_join_common_utils.cuh. See the developer guide on includes for more info.
Suggested change
| #include "mixed_join_kernel_semi_impl.cuh" | |
| #include "join/mixed_join_kernel_semi_impl.cuh" |
| */ | ||
|
|
||
| #include "join/mixed_join_common_utils.cuh" | ||
| #include "mixed_join_kernel_semi_impl.cuh" |
There was a problem hiding this comment.
Suggested change
| #include "mixed_join_kernel_semi_impl.cuh" | |
| #include "join/mixed_join_kernel_semi_impl.cuh" |
github-actions
bot
added
Python
wence-
removed
Python
tgujar
force-pushed
the
distinct-join-occupancy
branch
from
19944cf to
448b14b
Compare
|
Thanks for the review! Closing this PR, this work is merged into #15700. Please take a look there for further review |
3 tasks
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Description
Adds specialized dispatch for distinct hash join similar to that implemented for hash joins in #15700
Related issue: #15502
Checklist
Benchmark results on A100