Skip to content

Support concurrent aggregate calls on one streaming_groupby - #23884

Draft
PointKernel wants to merge 2 commits into
NVIDIA:mainfrom
PointKernel:streaming-groupby-concurrent-aggregate
Draft

Support concurrent aggregate calls on one streaming_groupby#23884
PointKernel wants to merge 2 commits into
NVIDIA:mainfrom
PointKernel:streaming-groupby-concurrent-aggregate

Conversation

@PointKernel

Copy link
Copy Markdown
Member

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_keys becomes std::atomic because concurrent aggregate() otherwise races the read in distinct_keys(). merge() shares probe_and_insert and takes the same lock. Insertion kernels from different batches still do not overlap, which #23428 allows for a first pass, and thrust::copy_if inside insertion already forces a host-visible sync.

Checklist

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

@copy-pr-bot

copy-pr-bot Bot commented Aug 29, 2026

Copy link
Copy Markdown

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.

@github-actions github-actions Bot added the libcudf Affects libcudf (C++/CUDA) code. label Aug 29, 2026
@PointKernel PointKernel added feature request New feature or request non-breaking Non-breaking change labels Aug 29, 2026
@PointKernel

Copy link
Copy Markdown
Member Author

/ok to test

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

feature request New feature or request libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[FEA] Support concurrent aggregate calls on one streaming_groupby instance

1 participant