[FEA]: Remove support for CDPv1 in Thrust algorithms

[jrhemstad]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

Thrust

Is your feature request related to a problem? Please describe.

As described in #1137, CDPv2 removed the ability to do device-side synchronization. As a result, Thrust algorithms invoked in device code with the thrust::device execution policy implicitly fall back to using thrust::seq.

However, Thrust preserved the ability to still use CDPv1 on architectures/CTKs where it is still supported. This has proven to be costly to preserve and deemed no longer worth the effort.

In other words, this kernel has different behavior depending on how it is compiled:

__global__ void kernel(){
   thrust::reduce(thrust::device, ...);
}

1. Launches a kernel (CDPv1)

   • Only if -rdc=true -arch=sm_N for N < 90 and CTK < 12.0

2. Falls back to thrust::seq

   • If -arch=sm_N for N >= 90 or CTK >= 12.0

Describe the solution you'd like

Invoking thrust::algorithm(thrust::device,...) in device code should always be equivalent to thrust::allgorithm(thrust::seq,...) (case 2 above).

Describe alternatives you've considered

No response

Additional context

No response