Use stream pool for gather/scatter. #14162
Open
Conversation
This file contains hidden or 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
|
Gather (scatter) summary:
|
|
This is probably not going to do much for the listed issue. The issue here is the number of raw thrust and kernel calls stemming from nesting (think thousands of columns underneath the top level table). The fix for that is going to be smarter parallelization of what is currently the recursive cpu-side approach. We have some ideas here. |
abellina
reviewed
Sep 29, 2023
| // only a single column, the fork/join overhead should be avoided. | ||
| auto streams = std::vector<rmm::cuda_stream_view>{}; | ||
| if (num_columns > 1) { | ||
| streams = cudf::detail::fork_streams(stream, num_columns); |
There was a problem hiding this comment.
I am curious how this works for per-thread default stream. For spark, we build cuDF with PTDS. Will streams will be have number-of-columns vector of the PTDS stream?
There was a problem hiding this comment.
The stream passed in would be the PTDS stream. Then a stream pool (for that thread) would be created (or reused), the work would be executed across that stream pool, and then the join step would insert events for all the elements of streams to be synchronized with stream before new work on stream (the PTDS stream in Spark's case) would be runnable.
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
bdice
commented
Mar 10, 2025
|
@mythrocks Could you pick this up? I just merged some changes that pull it up-to-date with |
I'm on it, mate. Thank you for this change. I'll post here with any findings. |
bdice
commented
Mar 10, 2025
vuule
reviewed
Mar 10, 2025
| // only a single column, the fork/join overhead should be avoided. | ||
| auto streams = std::vector<rmm::cuda_stream_view>{}; | ||
| if (num_columns > 1) { | ||
| streams = cudf::detail::fork_streams(stream, num_columns); |
There was a problem hiding this comment.
this will emit a warning when there are more columns than streams in the pool. We can remove with warning; otherwise it's probably good to limit the number of streams to the pool size.
There was a problem hiding this comment.
Is there a good/recommended number? And is it tuned by hand or automatically computed?
There was a problem hiding this comment.
It's fine to use all streams in the pool, I don't think we need to tune this.
but getting more reuses them in round-robin fashion and we have a warning because maybe it's an unexpected behavior for users.
There was a problem hiding this comment.
I knew we did round-robin, I did not know we issued a warning. I think removing the warning would be appropriate.
vuule
reviewed
Mar 10, 2025
| } | ||
|
|
||
| auto it = thrust::make_counting_iterator<size_type>(0); | ||
|
|
||
| std::transform(it, it + num_columns, result.begin(), [&](size_type i) { |
There was a problem hiding this comment.
thrust::tabulate saves a bit of code here
mythrocks
reviewed
Mar 11, 2025
Co-authored-by: MithunR <[email protected]>
ttnghia
approved these changes
Mar 17, 2025
|
I'll retarget this to 25.06. This will be good to have, but it didn't help NVIDIA/spark-rapids#12195. It simply wasn't the slow path, in this case. :/ |
3 tasks
|
Moved to 25.08. @mythrocks said that he should be able to get this done for the next release. |
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
This PR uses the stream pool introduced in #13922 to gather/scatter each column in a table on a separate stream.
Related: #13509, which this might resolve (need to verify).
Checklist