Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Avoid thrust logical operation. #9199

Merged
merged 2 commits into from
May 26, 2023
Merged

Conversation

trivialfis
Copy link
Member

@trivialfis trivialfis commented May 25, 2023

Thrust implementation of thrust::all_of/any_of/none_of adopts an early stopping strategy to bailout early by dividing the input into small batches. This is not ideal for data validation as we expect all data to be valid. The strategy leads to excessive kernel launches and stream synchronization.

A small optimization from #9194 .

Thrust implementation of `thrust::all_of/any_of/none_of` adopts an early stopping strategy
to bailout early by dividing input into small batches. This is not idea for data
validation as we expect all data to be valid. The strategy leads to excessive kernel
launches and stream synchronization.
@trivialfis
Copy link
Member Author

@rongou Would be great if you can take a look when you are available.

// intervals to early stop. But we expect all data to be valid here, using small
// intervals only decreases performance due to excessive kernel launch and stream
// synchronization.
auto valid = dh::Reduce(thrust::cuda::par(alloc), value_iter, value_iter + batch.Size(), true,
Copy link
Contributor

Choose a reason for hiding this comment

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

Would it be simpler to use thrust::transform_reduce?

Copy link
Member Author

Choose a reason for hiding this comment

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

We have trouble with using thrust reduce directly as it doesn't support input larger than INT_MAX.

Copy link
Contributor

Choose a reason for hiding this comment

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

Never mind, looks like we need the index to iterate.

@trivialfis trivialfis merged commit 053aaba into dmlc:master May 26, 2023
@trivialfis trivialfis deleted the prevent-thrust-es branch May 26, 2023 17:37
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants