[EPIC] Consolidate kernels between Thrust and CUB

[jrhemstad]

Below is a list of tasks in prioritized order. We should start with algorithms that already exist in CUB. This will allow delivering CUB optimizations into Thrust sooner.

Document procedure

• Author contributor guide for consolidating Thrust/CUB kernels/algorithms #1585

Replace Thrust Algorithms with CUB

• Refactor thrust::unique_by_key to use cub::DeviceSelect::UniqueByKey #1210
• Refactor thrust::copy_if to use cub::DeviceSelect #1263
• Refactor thrust::partition to use cub::DevicePartition #1383
• Refactor thrust::partition_copy and thrust::stable_partition_copy to use DevicePartition #1397
• Refactor thrust::unique to use cub::DeviceSelect::Unique #1625
• Refactor thrust::reduce to use cub::DeviceReduce #1626
• Refactor thrust::reduce_by_key to use cub::DeviceReduce::ReduceByKey
• Refactor thrust/extrema.h to use cub::DeviceReduce
• Make sure that thrust calls to cub algorithms use the Device* interface (i.e., to make sure we use the index type that is optimized for the cub algorithm) - or alternatively go via Dispatch* interface but make sure to use the right offset type

Port Thrust Algorithms into CUB

• Port thrust::cuda_cub::parallel_for to CUB #1231
• Port thrust::merge to CUB #1763
• Port thrust/thrust/system/cuda/detail/set_operations.h to CUB

A few notes:

• thrust::partition_copy and thrust::stable_partition_copy require taking two separate/distinct output iterators: one for the selected and one for the rejected items. DevicePartition, however, currently only supports a single output iterator, where the selected items are written to the beginning in order and the rejected items are written to the end in reverse order, respectively. Supporting the these two thrust algorithms requires extending AgentSelectIf, implementing overloads for methods like ScatterTwoPhase that are concerned with writing rejected items to the output iterators.
• Partition and copy doesn't have stencil overload in CUB. We might experiment with using ::Flagged version along with a transform iterator.
• While porting thrust::reduce_by_key we need to decide on accumulator type.

[gevtushenko]

Blocked by

• [EPIC] Implement VSMem abstraction #548