[BUG]: GH200 DeviceReduce performance: 14x (<1 GiB) and 2x (>1 GiB) lower than SOL

### Is this a duplicate?

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

### Type of Bug

Performance

### Component

Thrust

### Describe the bug

Improve the performance of `thrust::reduce` and `transform_reduce` by 14x for < 1 GB input sizes, and by 2x for >1 GiB sizes.

This requires fixing the following bugs:
- [x] https://github.com/NVIDIA/cccl/issues/435
- [ ] https://github.com/NVIDIA/cccl/issues/436
- [ ] Avoid stream sync in `DeviceReduce`.
- [ ] (only for non-associative reductions): https://github.com/NVIDIA/cccl/issues/265

### How to Reproduce

Run `DeviceReduce` and compare it against SOL throughput.

### Expected behavior

`DeviceReduce` should not be more than an order-of-magnitude slower than SOL.

### Reproduction link

Internal link available.

### Operating System

Linux.

### nvidia-smi output

GH200

### NVCC version

Any.