Skip to content

[FEA] Performance improvement for mixed left semi/anti join - #15288

Merged
rapids-bot[bot] merged 8 commits into
NVIDIA:branch-24.06from
tgujar:mixed_semi_join_perf
Apr 3, 2024
Merged

[FEA] Performance improvement for mixed left semi/anti join#15288
rapids-bot[bot] merged 8 commits into
NVIDIA:branch-24.06from
tgujar:mixed_semi_join_perf

Conversation

@tgujar

@tgujar tgujar commented Mar 13, 2024

Copy link
Copy Markdown
Contributor

Current implementation of mixed semi/anti join probes the built hash table twice -- once to find the output table size and once to build the output. Since the upper bound on output table size is O(N) where N is the size of the left table, we can avoid probing twice and achieve a faster join implementation.

This implementation reserves the required upper memory bound, builds the output, and then collects the relevant output rows. This probes the hash table only once.

This PR also removes the size kernels for mixed semi join and output size parameters passed to the mixed semi join.

Closes #15250

Benchmark Results from cudf repository

mixed_left_semi_join_32bit (New implementation)

[0] NVIDIA TITAN V

| Key Type | Payload Type | Nullable | Build Table Size | Probe Table Size | Samples |  CPU Time  | Noise |  GPU Time  | Noise |
|----------|--------------|----------|------------------|------------------|---------|------------|-------|------------|-------|
|      I32 |          I32 |        0 |           100000 |           100000 |   1920x | 266.239 us | 3.43% | 261.324 us | 2.84% |
|      I32 |          I32 |        0 |           100000 |           400000 |   1024x | 495.434 us | 1.18% | 490.544 us | 0.63% |
|      I32 |          I32 |        0 |         10000000 |         10000000 |     24x |  20.919 ms | 0.04% |  20.914 ms | 0.03% |
|      I32 |          I32 |        0 |         10000000 |         40000000 |     11x |  54.697 ms | 0.03% |  54.692 ms | 0.03% |
|      I32 |          I32 |        0 |         10000000 |        100000000 |     11x | 122.171 ms | 0.03% | 122.166 ms | 0.03% |
|      I32 |          I32 |        0 |         80000000 |        100000000 |     11x | 192.979 ms | 0.01% | 192.975 ms | 0.01% |
|      I32 |          I32 |        0 |        100000000 |        100000000 |     11x | 212.878 ms | 0.01% | 212.874 ms | 0.01% |
|      I32 |          I32 |        0 |         10000000 |        240000000 |     11x | 279.794 ms | 0.01% | 279.790 ms | 0.01% |
|      I32 |          I32 |        0 |         80000000 |        240000000 |     11x | 351.186 ms | 0.01% | 351.183 ms | 0.01% |
|      I32 |          I32 |        0 |        100000000 |        240000000 |     11x | 370.794 ms | 0.01% | 370.790 ms | 0.01% |

mixed_left_semi_join_32bit (Old implementation)

[0] NVIDIA TITAN V

| Key Type | Payload Type | Nullable | Build Table Size | Probe Table Size | Samples |  CPU Time  | Noise |  GPU Time  | Noise |
|----------|--------------|----------|------------------|------------------|---------|------------|-------|------------|-------|
|      I32 |          I32 |        0 |           100000 |           100000 |   1392x | 368.030 us | 3.05% | 363.065 us | 2.70% |
|      I32 |          I32 |        0 |           100000 |           400000 |    832x | 832.492 us | 0.84% | 827.586 us | 0.60% |
|      I32 |          I32 |        0 |         10000000 |         10000000 |     16x |  32.310 ms | 0.03% |  32.305 ms | 0.03% |
|      I32 |          I32 |        0 |         10000000 |         40000000 |     11x | 100.222 ms | 0.03% | 100.218 ms | 0.03% |
|      I32 |          I32 |        0 |         10000000 |        100000000 |     11x | 235.874 ms | 0.01% | 235.870 ms | 0.01% |
|      I32 |          I32 |        0 |         80000000 |        100000000 |     11x | 307.042 ms | 0.01% | 307.038 ms | 0.01% |
|      I32 |          I32 |        0 |        100000000 |        100000000 |     11x | 326.797 ms | 0.01% | 326.794 ms | 0.01% |
|      I32 |          I32 |        0 |         10000000 |        240000000 |     11x | 552.730 ms | 0.01% | 552.728 ms | 0.01% |
|      I32 |          I32 |        0 |         80000000 |        240000000 |     11x | 624.958 ms | 0.01% | 624.956 ms | 0.01% |
|      I32 |          I32 |        0 |        100000000 |        240000000 |     11x | 644.148 ms | 0.00% | 644.146 ms | 0.00% |

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

Current implementation of mixed semi/anti join probes the built hash
table twice -- once to find the output table size and once to build the
output. Since the upper bound on output table size is O(N) where N is
the size of the left table, we can avoid probing twice and achieve a
faster join implementation.

This implementation reserves the required upper memory bound, builds the
output, and then collects the relevant output rows. This probes the hash
table only once.

- Closes NVIDIA#15250
@tgujar
tgujar requested review from a team as code owners March 13, 2024 16:39
@tgujar
tgujar requested review from bdice and mhaseeb123 March 13, 2024 16:39
@copy-pr-bot

copy-pr-bot Bot commented Mar 13, 2024

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@github-actions github-actions Bot added libcudf Affects libcudf (C++/CUDA) code. CMake CMake build issue labels Mar 13, 2024
@tgujar tgujar changed the title Performance improvement for mixed left semi/anti join [FEA] Performance improvement for mixed left semi/anti join Mar 13, 2024

@bdice bdice left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nice. Overall, this is close to what I was hoping for! I have one algorithmic suggestion that would reduce the lines of code. I haven't benchmarked or built my suggestion, though, so feel free to re-measure (I expect compile time and binary size to decrease, but runtime might be a tiny bit slower?).

Comment thread cpp/src/join/mixed_join_semi.cu

@mhaseeb123 mhaseeb123 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hi @tgujar, I am new to libcudf team (joined just last Monday) but I have gone through the code changes as thoroughly as possible. I mostly understood your new code and I like the changes and the performance improvements so far. I saw @bdice's suggestion on building gather_map as boolean_mask(thrust::sequence, predicate) instead of with copy_if and I would also like to see if it could be done that way and if there are any performance consequences to it. Thank you!

@PointKernel PointKernel added Performance Performance related issue improvement Improvement / enhancement to an existing function non-breaking Non-breaking change breaking Breaking change and removed non-breaking Non-breaking change labels Mar 13, 2024
Comment thread cpp/include/cudf/join.hpp
Comment thread cpp/src/join/mixed_join_semi.cu
Comment thread cpp/src/join/mixed_join_semi.cu Outdated
Comment thread cpp/include/cudf/join.hpp
Comment thread cpp/src/join/mixed_join_semi.cu Outdated
Comment thread cpp/src/join/mixed_join_semi.cu
Comment on lines +259 to +267
// gather_map_end will be the end of valid data in gather_map
auto gather_map_end =
thrust::copy_if(rmm::exec_policy(stream),
thrust::counting_iterator<size_type>(0),
thrust::counting_iterator<size_type>(probe.num_rows()),
gather_map->begin(),
[join_type, d_flagged = flagged->begin()] __device__(size_type const idx) {
return *(d_flagged + idx) == (join_type == detail::join_kind::LEFT_SEMI_JOIN);
});

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

suggestion which may not be better, rather than closing over the flagged iterator in the lambda, how about using the stencil form of copy_if

auto gather_map_end =
  thrust::copy_if(rmm::exec_policy(stream),
                  thrust::counting_iterator<size_type>(0),
                  thrust::counting_iterator<size_type>(probe.num_rows()),
                  flagged.begin(),
                  gather_map->begin(),
                  [join_type] __device__(bool flag) {
                    return flag == (join_type == detail::join_kind::LEFT_SEMI_JOIN);
                  });

WDYT?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

fixed, I think this is better, thanks!

Comment thread cpp/include/cudf/join.hpp
* be relied upon, simply passed to the corresponding `mixed_left_join` API as
* is.
*/
std::pair<std::size_t, std::unique_ptr<rmm::device_uvector<size_type>>> mixed_left_semi_join_size(

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

question: (cc @revans2) Does spark use these APIs for estimating whether a join must be batched because it would overflow memory?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I want to note here that this function itself reserves memory equivalent to the max output size of the semi join.

@wence-

wence- commented Mar 14, 2024

Copy link
Copy Markdown
Contributor

/ok to test

tgujar and others added 2 commits March 14, 2024 07:57
- Use default resource for scratch mem allocation.
- Use stencil version of copy_if
@PointKernel

Copy link
Copy Markdown
Member

/ok to test

Comment thread cpp/src/join/mixed_join_kernels_semi.cu Outdated
Comment thread cpp/src/join/mixed_join_kernels_semi.cu Outdated
Comment thread cpp/src/join/mixed_join_semi.cu
Comment thread cpp/src/join/mixed_join_semi.cu Outdated
tgujar added 2 commits March 18, 2024 08:55
Initialization kernel can be skipped if we write to flagged array on
non-match in hash table.
Remove function calls for mixed semi/anti join size calculation
and remove optional output size parameter to conform with changes to
API in libcudf
@tgujar
tgujar requested a review from a team as a code owner March 19, 2024 22:37
@github-actions github-actions Bot added the Java Affects Java cuDF API. label Mar 19, 2024
@jlowe

jlowe commented Mar 20, 2024

Copy link
Copy Markdown
Contributor

/ok to test

@jlowe jlowe left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Java approval

@PointKernel

Copy link
Copy Markdown
Member

/ok to test

@PointKernel PointKernel left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Great work! Thanks for helping improve cudf!

@wence-

wence- commented Mar 20, 2024

Copy link
Copy Markdown
Contributor

As a followup, which would be an API-break and a change from the normal (not semi-) joins, I wonder if it is worthwhile removing the thrust::copy_if that converts the boolean keep mask into a gather map. We could instead just return the mask directly and then compute the output of the merge with apply_boolean_mask rather than gather. This would reduce the memory footprint (since without computing the number of true values in the keep mask, we must first allocate space for num_rows integers and then compact after the thrust::copy_if call), and maybe perform better. However, it may just move the costs around a bit, so not necessarily an improvement.

@tgujar

tgujar commented Mar 20, 2024

Copy link
Copy Markdown
Contributor Author

As a followup, which would be an API-break and a change from the normal (not semi-) joins, I wonder if it is worthwhile removing the thrust::copy_if that converts the boolean keep mask into a gather map. We could instead just return the mask directly and then compute the output of the merge with apply_boolean_mask rather than gather. This would reduce the memory footprint (since without computing the number of true values in the keep mask, we must first allocate space for num_rows integers and then compact after the thrust::copy_if call), and maybe perform better. However, it may just move the costs around a bit, so not necessarily an improvement.

If this table is involved in another join operation, we would have to first materialize the rows numbers if I understand correctly. I am not sure if the performance difference would be worth the additional complexity.

For memory footprint, we could add a prefix sum kernel and only reserve the required memory, but this may not be a good tradeoff

@wence-

wence- commented Mar 21, 2024

Copy link
Copy Markdown
Contributor

As a followup, which would be an API-break and a change from the normal (not semi-) joins, I wonder if it is worthwhile removing the thrust::copy_if that converts the boolean keep mask into a gather map. We could instead just return the mask directly and then compute the output of the merge with apply_boolean_mask rather than gather. This would reduce the memory footprint (since without computing the number of true values in the keep mask, we must first allocate space for num_rows integers and then compact after the thrust::copy_if call), and maybe perform better. However, it may just move the costs around a bit, so not necessarily an improvement.

If this table is involved in another join operation, we would have to first materialize the rows numbers if I understand correctly. I am not sure if the performance difference would be worth the additional complexity.

For memory footprint, we could add a prefix sum kernel and only reserve the required memory, but this may not be a good tradeoff

Well, the gather map that is currently returned is just the sparse representation of the boolean mask. So you would move from doing:

gather_map = semi_join(left, right);
result = cudf::gather(left, gather_map);

To:

mask = semi_join(left, right);
result = cudf::apply_boolean_mask(left, mask);

That puts (in many cases) the responsibility of constructing the sparse gather map on apply_boolean_mask (which does run a prefix-sum kernel to not over-allocate), but devolves to thrust::copy_if directly from the mask into the output in cases where it is possible.

@tgujar
tgujar changed the base branch from branch-24.04 to branch-24.06 March 23, 2024 19:17
@PointKernel

Copy link
Copy Markdown
Member

/ok to test

@PointKernel

Copy link
Copy Markdown
Member

/merge

@rapids-bot
rapids-bot Bot merged commit fbaad8a into NVIDIA:branch-24.06 Apr 3, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

breaking Breaking change CMake CMake build issue improvement Improvement / enhancement to an existing function Java Affects Java cuDF API. libcudf Affects libcudf (C++/CUDA) code. Performance Performance related issue

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[FEA] Performance improvement for mixed semi joins

7 participants