-
-
Notifications
You must be signed in to change notification settings - Fork 8.7k
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
Enable compiling with system cub. #7232
Conversation
Eventually, we will have to move on to using system cub, with ctk 11.4 and custom cub there are lots of warnings. I think this might be the last version we can workaround it. |
be6dbc0
to
ab7d37d
Compare
}); | ||
// shrink down to pair | ||
auto fptp_it_out = thrust::make_transform_output_iterator( | ||
dh::tbegin(d_fptp), [=] __device__(Triple const &t) { | ||
return thrust::make_pair(thrust::get<1>(t), thrust::get<2>(t)); | ||
dh::TypedDiscard<Triple>{}, [d_fptp] __device__(Triple const &t) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Originally this reduces the triple to pair, cub in CUDA 11.0~11.2 doesn't handle changed type. So here I just write the result in lambda and use discard iterator to keep the output type unchanged.
* For now we need to limit the number of items to INT32_MAX This is caused by latest RMM brings cub into its include path, so when XGBoost is compiled with RMM the paths conflict.
ab7d37d
to
5a3a1ae
Compare
@trivialfis Could you cherry-pick this PR to 1.4.0 branch? |
No, I'm planning for a new release |
Enable compiling with system cub. (dmlc#7232) See merge request nvspark/xgboost!394
- This works only with CUDA 11.4, for CUDA 11.2 there are some obscure errors with inclusive scan and custom iterators.