statically detect precondition violations in thrust::for_each

Passing an `InputIterator` to `thrust::for_each` should probably be a precondition violation, which ought to be detected statically, e.g., with a `static_assert`.

Currently, when passing such an iterator, what happens is the following:
- `thrust::distance` is called [here](https://github.com/NVIDIA/thrust/blob/24486a169a62a58ef8f824d3dc9613c006b6f5a7/thrust/system/cuda/detail/for_each.h#L98),
- then `cuda_cub::for_each_n` is called with an "exhausted" input iterator ([here](https://github.com/NVIDIA/thrust/blob/24486a169a62a58ef8f824d3dc9613c006b6f5a7/thrust/system/cuda/detail/for_each.h#L99)),
- then `cuda_cub::parallel_for` is called [here](https://github.com/NVIDIA/thrust/blob/24486a169a62a58ef8f824d3dc9613c006b6f5a7/thrust/system/cuda/detail/for_each.h#L80-L82), which launches a kernel that copies the exhausted input iterator to all threads, and a length of the number of exhausted elements, and
- finally the kernel tries to advance each exhausted input iterator [here](https://github.com/NVIDIA/thrust/blob/24486a169a62a58ef8f824d3dc9613c006b6f5a7/thrust/system/cuda/detail/parallel_for.h#L97), by using `*(it + i)` [here](https://github.com/NVIDIA/thrust/blob/24486a169a62a58ef8f824d3dc9613c006b6f5a7/thrust/system/cuda/detail/for_each.h#L58), which may compile if the input iterator supports advancing by `it + i`.

I think that:
- thrust::for_each should check that the iterators are forward iterators early, and static assert if they aren't.
- the cuda_cub::parallel_for dispatch path that `*(it + i)` should only be used if the iterators are random access iterators
- if they aren't either static assert, or
- copy the forward iterators to each device thread, and use `thrust::advance` to advance them in parallel