Gather inventory of CuPy's CCCL usage

[jrhemstad]

We'd like a list of:

• Every Thrust/CUB function used in CuPy today
• Thrust/CUB functions CuPy would like to be using but cannot

[shwina]

Every Thrust/CUB function used in CuPy today

This is just an initial analysis. CUB and Thrust usage in CuPy can be found in cupy_cub.cu and cupy_thrust.cu respectively.

cub

• DeviceReduce
• DeviceSegmentedReduce
• DeviceScan
• DeviceSpmv
• DeviceHistogram

thrust

• stable_sort
• stable_sort_by_key

For sorting a muilti-dimensional array along a specified axis (dimension), there’s a specialized algorithm which additionally uses:

• sequence
• transform
• thrust::device_pointer_cast
• thrust::device_ptr
• thrust::tuple, thrust::get
• thrust::make_counting_iterator
• thrust::make_constant_iterator
• thrust::divides

cub “patching”

In addition, within cupy_cub.cu there’s some code that specializes/supplements CUB:

• Specializations like numeric_limits<thrust::complex<float>>> to support complex numbers
• Specializing operators like Max(), Min() etc., for floats to handle NaNs in a way that matches NumPy semantics.

I’m not entirely sure how to include the the above specializations when switching CuPy over to use cuda.parallel. Something to keep in mind.

[leofang]

This is just an initial analysis. CUB and Thrust usage in CuPy can be found in cupy_cub.cu and cupy_thrust.cu respectively.

The list looks fine to me.

[leofang]

I suggested during the meeting that we deprioritize DeviceHistogram because of several issues, e.g. #912