Support concurrent aggregate calls on one streaming_groupby - #23884
Draft
PointKernel wants to merge 2 commits into
Draft
Support concurrent aggregate calls on one streaming_groupby#23884PointKernel wants to merge 2 commits into
PointKernel wants to merge 2 commits into
Conversation
|
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. |
Member
Author
|
/ok to test |
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
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
Closes #23428
This PR makes
streaming_groupby::aggregate()safe to call concurrently from multiple host threads on one instance, each with its own stream, without caller-side serialization. Newly discovered keys live in the hash set under a transient encoding,max_distinct_keys + row_idx, that carries no batch identifier, so overlapping insertions decode each other's values against the wrong batch table. The insertion phase also mutates shared host state that must stay in lockstep: the batch ID, the dense ID base, and the retained key batches.Rather than widen the encoding to carry a batch ID, a mutex serializes the insertion phase and a CUDA event orders it across calls on different streams. The aggregation that follows stays outside the lock, since it updates every group through
cudf::detail::atomic_add/atomic_min/atomic_max._distinct_keysbecomesstd::atomicbecause concurrentaggregate()otherwise races the read indistinct_keys().merge()sharesprobe_and_insertand takes the same lock. Insertion kernels from different batches still do not overlap, which #23428 allows for a first pass, andthrust::copy_ifinside insertion already forces a host-visible sync.Checklist