thrust::*_into algorithms to expose more asynchrony

We should have `thrust::reduce_into(first, last, out)` that is conceptually equivalent to `*out = thrust::reduce(first, last)`. That would mean `reduce_into` is no longer data dependent, so if you pass par_nosync it won't have to block. This is also useful for avoiding a transfer from device to host if you're just going to use the reduction result on the device.

All the data dependent Thrust algorithms should have an `_into` form. That means `reduce`, `count`, `find`, etc.