Skip to content

Optimize left_semi_join by materializing the gather mask - #10511

Merged
rapids-bot[bot] merged 3 commits into
NVIDIA:branch-22.06from
cheinger:branch-22.06
May 5, 2022
Merged

Optimize left_semi_join by materializing the gather mask#10511
rapids-bot[bot] merged 3 commits into
NVIDIA:branch-22.06from
cheinger:branch-22.06

Conversation

@cheinger

@cheinger cheinger commented Mar 24, 2022 •
edited by jrhemstad
Loading

Copy link
Copy Markdown
Contributor

Closes #10464

Updates the left_semi_join to materialize the gather mask instead of generating it via a transform iterator.

Including the map.contains in the gather call reduced occupancy due to increasing register usage. As a result, explicitly materializing the gather mask is faster.

@cheinger
cheinger requested a review from a team as a code owner March 24, 2022 20:02
@cheinger
cheinger requested review from codereport and trxcllnt March 24, 2022 20:02
@GPUtester

GPUtester commented Mar 24, 2022

Copy link
Copy Markdown
Collaborator

Can one of the admins verify this patch?

@github-actions github-actions Bot added the libcudf Affects libcudf (C++/CUDA) code. label Mar 24, 2022
@cheinger cheinger mentioned this pull request Mar 24, 2022
Closed
jrhemstad
jrhemstad reviewed Mar 24, 2022
View reviewed changes
Comment thread cpp/src/join/semi_join.cu Outdated
jrhemstad
jrhemstad reviewed Mar 24, 2022
View reviewed changes
Comment thread cpp/src/join/semi_join.cu Outdated
jrhemstad
jrhemstad reviewed Mar 24, 2022
View reviewed changes
Comment thread cpp/src/join/semi_join.cu Outdated
jrhemstad
jrhemstad reviewed Mar 24, 2022
View reviewed changes
Comment thread cpp/src/join/semi_join.cu Outdated
jrhemstad
jrhemstad reviewed Mar 24, 2022
View reviewed changes
Comment thread cpp/src/join/semi_join.cu Outdated
@PointKernel PointKernel added improvement Improvement / enhancement to an existing function non-breaking Non-breaking change labels Mar 29, 2022
@xavier-nvidia
Optimized left_semi_join
52df19b 
Up to 20x faster. Separated hash table lookup from copy_if because
increased register usage significantly limited occupancy of this
kernel.
@jrhemstad

jrhemstad commented Apr 5, 2022 •
edited
Loading

Copy link
Copy Markdown
Contributor

@cheinger so to be clear, the performance improvement didn't come from using cub::DeviceSelect::Flagged, but instead from pulling the map::contains function out of the copy_if by materializing the predicate as a separate array?

@cheinger

cheinger commented Apr 5, 2022

Copy link
Copy Markdown
Contributor Author

@jrhemstad correct. I updated the gitlab issue with a more detailed explanation

@jrhemstad

jrhemstad commented Apr 5, 2022

Copy link
Copy Markdown
Contributor

@cheinger could you update the PR description to provide a short summary? The PR description goes into the CHANGELOG.

@PointKernel

PointKernel commented Apr 5, 2022

Copy link
Copy Markdown
Member

ok to test

PointKernel
PointKernel reviewed Apr 5, 2022
View reviewed changes

@PointKernel PointKernel left a comment •
edited
Loading

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.

Nice work. Can you please update the PR title accordingly? It would be useful to also include your performance analysis (here) in the PR description. Did you notice any performance changes in semi join benchmarks?

Comment thread cpp/src/join/semi_join.cu
Comment thread cpp/src/join/semi_join.cu Outdated
@PointKernel PointKernel added the Performance Performance related issue label Apr 5, 2022
@sevagh

sevagh commented Apr 5, 2022

Copy link
Copy Markdown
Contributor

ok to test

1 similar comment
@sevagh

sevagh commented Apr 5, 2022

Copy link
Copy Markdown
Contributor

ok to test

@codecov

This comment was marked as outdated.

Sign in to view
@jrhemstad

jrhemstad commented May 5, 2022

Copy link
Copy Markdown
Contributor

@PointKernel can you re-review/approve?

@PointKernel PointKernel changed the title Optimize left_semi_join by using cub::DeviceSelect::Flagged instead of thrust::copy_if Optimize left_semi_join by materializing the gather mask May 5, 2022
@jrhemstad

jrhemstad commented May 5, 2022

Copy link
Copy Markdown
Contributor

add to whitelist

@ajschmidt8

ajschmidt8 commented May 5, 2022

Copy link
Copy Markdown
Member

add to allowlist

@PointKernel

PointKernel commented May 5, 2022

Copy link
Copy Markdown
Member

rerun tests

@PointKernel

PointKernel commented May 5, 2022

Copy link
Copy Markdown
Member

@gpucibot merge

@rapids-bot
rapids-bot Bot merged commit ee26fbe into NVIDIA:branch-22.06 May 5, 2022
@jrhemstad jrhemstad mentioned this pull request May 26, 2022
Merged
@GregoryKimball

GregoryKimball commented May 31, 2022

Copy link
Copy Markdown
Collaborator

Thank you @cheinger for adding this optimization! I'm seeing a 15-30% reduction in compute time for our JOIN benchmarks as a result of this change.
image

@cheinger

cheinger commented May 31, 2022

Copy link
Copy Markdown
Contributor Author

@GregoryKimball Sweet! Happy to help!

@ttnghia ttnghia mentioned this pull request Jun 23, 2022
Merged
@ttnghia ttnghia mentioned this pull request Jul 6, 2022
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@PointKernel PointKernel PointKernel approved these changes
@jrhemstad jrhemstad jrhemstad approved these changes
@trxcllnt trxcllnt Awaiting requested review from trxcllnt trxcllnt was automatically assigned from rapidsai/cudf-cpp-codeowners
@codereport codereport Awaiting requested review from codereport codereport was automatically assigned from rapidsai/cudf-cpp-codeowners

Assignees

No one assigned

Labels

improvement Improvement / enhancement to an existing function libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change Performance Performance related issue

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

[BUG] Left Semi Join much slower than Inner join

9 participants

@cheinger @GPUtester @jrhemstad @PointKernel @sevagh @ajschmidt8 @GregoryKimball @caryr35 @xavier-nvidia