[FEA]: Avoiding unnecessary work in BlockHistogram

### 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

CUB

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

I would like to be able to use `cub::BlockHistogram` to efficiently run-length-encode

1. pre-sorted data (avoiding internal sort) and/or
2. data that is expected to fill almost all bins with at least one value (avoiding internal scan).

### Describe the solution you'd like

1. Avoiding the internal sort in `BLOCK_HISTO_SORT` with a new algorithm - e.g. `BLOCK_HISTO_PRESORTED` is trivial to implement on its own, I see no reason not to do so.

2. This case is less straightforward. My idea is to avoid the scan at the end of `BLOCK_HISTO_SORT` and instead write out the end-indices of intermediate bins on the fly in the internal `DiscontinuityOp`, i.e. use a single `temp_storage.discontinuities.run_offset` in shared storage instead of `run_begin` and `run_end` and then replace https://github.com/NVIDIA/cccl/blob/05fadcfecef93a1aaed0c9165aed290aebad23fb/cub/cub/block/specializations/block_histogram_sort.cuh#L158-L159 with a loop writing out run-ends (`b_index`) in the range [`a + 1`, `min(b, BINS)`].

Assuming that the loop will almost always run only a single iteration due to the distribution of the input data, I would think that it should be faster than the current post-processing scan even though one might still need some post-processing to fill offsets for empty bins at the start and end of the histogram.

It might be possible to avoid initializing the offsets - and therefore synchronizing the block - using the sorted data, but these are details that would have to be figured out.

I see the slight problem that these two ideas are orthogonal, i.e. when implementing both one would end up with 3 new strategies
- without sort,
- without scan and
- without sort and without scan

### Describe alternatives you've considered

Right now I'm just using `cub::BlockDiscontinuity` directly in my code.

### Additional context

_No response_