[FEA]: Using `thrust::device` execution policy in device code should fail to compile

### Is this a duplicate?

- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Area

Thrust

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

As a user of Thrust, I might want to use a Thrust algorithm within the body of a kernel. Whether intentional or accidental, I might try and use the `thrust::device` execution policy. For example,

```
__global__
void kernel(...){
auto result = thrust::reduce(thrust::device, ....);
}
```

What should I expect actually happens here?

#### Context: CUDA Dynamic Parallelism (CDP)

The answer is complicated and relates to breaking changes that happened with CUDA Dynamic Parallelism (CDP) with the introduction of sm_90.

In short, prior to sm_90, you could do the following:
```
__global__ void CDPv1_kernel(){
kernel<<<...>>>(...); // device side launch, aka CDP
cudaDeviceSynchronize(); // This is the important piece
}
```
In other words, you could launch a child kernel, and then synchronize to wait for the results from that kernel. This is know as "CDPv1".

CDPv1 was [deprecated in CUDA 11.6 ](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-parallelism-synchronization) and removed in CUDA 12.0.

You can still launch child kernels, but you can no longer block and synchronize to wait for their results. This is known as **CDPv2**. The recommended practice instead of synchronizing is to launch an additional tail child kernel that consumes the result of the first child kernel.

```
__global__ void CDPv1_kernel(){
kernel<<<...>>>(result,...); // device side launch, aka CDP
// cudaDeviceSynchronize(); // Can't do this
tail_kernel<<<...>>>(result...) // need to do this
}
```

#### CDPv1, CDPv2 and Thrust

To understand how this relates to Thrust, one must first understand that Thrust algorithms are inherently synchronous. Consider `thrust::reduce` that returns its result. It needs to synchronize the work of the kernel in order to ensure the final result is available.

In CDPv1, using `thrust::reduce(thrust::device,...` in a kernel would launch a child kernel and synchronize before returning the result.

With CDPv2, this is no longer possible. Therefore, it is not possible to preserve Thrust algorithm's synchronous semantics with CDPv2.

This would mean any existing code using `thrust::device` in device code would break. To avoid breaking that code, a decision was made to make `thrust::device` implicitly fallback to `thrust::seq`, i.e., it no longer launches a child kernel and instead just runs sequentially on the calling thread.

We have since realized that this was a mistake for two main reasons:
1. It is too astonishing to users
2. It is too onerous to maintain

### Describe the solution you'd like

Using a Thrust algorithm with the `thrust::device` execution policy should fail to compile and emit a helpful diagnostic explaining why and directing the user to use `thrust::seq` instead if that's the behavior they want.

### Describe alternatives you've considered

An alternative would be to develop a more asynchronous programming model for Thrust where algorithms don't need to synchronize. `thrust::async` was a little known and undocumented attempt at doing this, but was quickly abandoned and we have no plans to revive it.

Furthermore, some algorithms don't _really_ need to synchronize because they don't return a result. This is part of the reason why we added the [`thrust::cuda::par_nosync`](https://github.com/NVIDIA/cccl/blob/0ff0f61c02c38c2ca25b92db06ff49a75140ca1f/thrust/examples/cuda/explicit_cuda_stream.cu#L40-L50) execution policy. But this execution policy is merely a _hint_. Some algorithms still require synchronization to be functionally correct.

We could explore adding a new execution policy similar to `thrust::cuda::par_nosync` that is only allowed on algorithms that truly do not need to synchronize. This could then be used for device-side usage of Thrust algorithms where launching a child kernel is desired.

However, all of this is purely hypothetical and we currently have no plans to explore it. There would need to be considerable demands from users for this kind of functionality for us to seriously consider investing time into exploring this.

### Additional context

For the record, CUB algorithms _do_ work with CDPv2 because CUB's programming model is asynchronous and more akin to launching a kernel directly. If someone truly wanted device-side kernel launch, then using an equivalent CUB algorithm would be the best way to do it.