Make sure `cub::DeviceScan` invokes the reduction operator only with valid data

_This issue is a sub-task of https://github.com/NVIDIA/cccl/issues/459. It focuses specifically on the `cub::DeviceScan` algorithm._

### Issue
As a user of Thrust/CUB, I can provide a custom reduction operator to the `cub::DeviceScan` algorithm. This operator may rely on some invariant I have on the provided data and, hence, only be well-defined for these values. As a result, it can be problematic if that operator is invoked on any other values. These invalid values can either be uninitialized values or repeated values, which means applying the reduction operator on an input item, `x_i` and a partial aggregate `agg`, op(`agg`, `x_i`), where `agg` already incorporated `x_i`.

### List of Related Issues
- https://github.com/NVIDIA/cccl/issues/780
- https://github.com/NVIDIA/cccl/issues/458
- https://github.com/NVIDIA/cccl/issues/801

### Potential Root Causes

For `cub::DeviceScan`, specifically, there may be multiple root causes:
1. The CUDA algorithm implementations within CUB will have threads in a block do things like default construct objects, read past boundaries even if they don't ultimately participate in the work, or repeatedly apply the reduction operation more than once on some items. This is _fine_ for simple operators like `+` but for custom operators it can cause failures.
2. The decoupled look-back algorithm processes windows of 32 tile states at a time. It waits until all tile states of this window are not flagged as `SCAN_TILE_INVALID` anymore. I.e., each tile state of that window contains one of the following three values: (1) a tile's aggregate, (2) a tile's inclusive prefix, or (3) an out-of-bounds value, which is some "uninitialized" value. The decoupled look-back algorithm will then process this window, computing a `TailSegmentedReduce` over this window to compute the inclusive prefix. `TailSegmentedReduce` starts from the "rightmost" tile for which the inclusive prefix is available computing the reduction operation over this window. If there's no inclusive prefix amongst any of the tile states, it will just compute the aggregate across this window and continue processing the next window. We need to investigate whether `TailSegmentedReduce` is mathematically correct, or whether the reduction operation is applied on some invalid data (e.g., invoking the reduction operation on an aggregate and an item `x`, where `x` was already incorporated in the aggregate, hence breaking the invariant the user has on their data).
3. The decoupled look-back algorithm initializes the the first 32 tile states with the `SCAN_TILE_OOB` flag, to make sure we can process the full window that includes the very first "valid" tile state without having to account for reading beyond the beginning of tile states. These out-of-bounds tiles states are processed by `TailSegmentedReduce` and, hence, are passing invalid data to the reduction operation.

### Tasks
- [x] https://github.com/NVIDIA/cccl/issues/5018
- [x] https://github.com/NVIDIA/cccl/issues/5019
- [ ] Mitigate any performance regressions, or expose the mathematically correct variant via the functional requirements API ([#1326](https://github.com/NVIDIA/cccl/pull/1326))