Many years ago we switched thrust::reduce/thrust::reduce_by_key to use CUB directly instead of using Thrust's old copy of CUB. This was done because of perf issues in the old copy of CUB reduce.
However, the old copy of CUB (which can be found here) was kept around, because some other algorithms used it. The hope was that we'd move those algorithms over next, then delete the deadcode. That doesn't seem to have happened.
I tried deleting this code, and it looks like at least the extrema functions (min_element, max_element, and minmax_element) use it. I think we could probably just not specialize these functions for the CUDA backend of Thrust, and instead have them go through the generic layer that will implement them in terms of reduce. However, CUB does have specific variants for min/max (but not minmax), so we may want to hook that up.
I'm not sure if those CUB variants are for perf or not - CUB should recognize the thrust::max/thrust::min function objects and dispatch to the fastpath.
Many years ago we switched
thrust::reduce/thrust::reduce_by_keyto use CUB directly instead of using Thrust's old copy of CUB. This was done because of perf issues in the old copy of CUB reduce.However, the old copy of CUB (which can be found here) was kept around, because some other algorithms used it. The hope was that we'd move those algorithms over next, then delete the deadcode. That doesn't seem to have happened.
I tried deleting this code, and it looks like at least the extrema functions (
min_element,max_element, andminmax_element) use it. I think we could probably just not specialize these functions for the CUDA backend of Thrust, and instead have them go through the generic layer that will implement them in terms of reduce. However, CUB does have specific variants for min/max (but not minmax), so we may want to hook that up.I'm not sure if those CUB variants are for perf or not - CUB should recognize the
thrust::max/thrust::minfunction objects and dispatch to the fastpath.