[BUG]: cub::DeviceHistogram::HistogramEven silently bins everything into bin 0 (or spuriously rejects the histogram) when the level range does not fit the bin computation's intermediate types

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

Silent Failure

### Component

Thrust

### Describe the bug

For integral sample/level types, `HistogramEven` misbehaves whenever `upper_level - lower_level` exceeds the sample type's maximum value — a perfectly ordinary request, e.g. histogramming the full value range of a type. Depending on the type width, the same request produces three different behaviors, none of them correct:

- **`int8_t` / `int16_t`: every sample is counted in bin 0 and `cudaSuccess` is returned.** A full-range `int16_t` histogram (levels [-32768, 32767)) with 100 bins reports success and puts all samples in bin 0 (verified). Same for `int8_t` over [-128, 127). There is no error, no NaN, no out-of-bounds — just a plausible-looking histogram that is wrong, which makes this easy to ship without noticing.
- **`int32_t`: the call is rejected with `cudaErrorInvalidValue`.** A histogram over [-1.5e9, 1.5e9) with 100 bins (verified) is refused even though nothing about it is invalid: the documentation only promises rejection when `(upper_level - lower_level) * (num_levels - 1)` exceeds `uint64_t`, and this product (3e11) is nowhere near that.
- **`int64_t`: such spans are also rejected**, which for this width is at least consistent with the documented overflow condition — but the rejection is reached through signed-overflow UB rather than by a valid check, so it is correct only by accident.

A related inconsistency affects the *number of bins*: bin counts exceeding what the level type can represent are supposed to be rejected (behavior introduced by #6908, and covered by a test), but the rejection only fires for some bin counts. Others slip through and silently compute a histogram with a *different bin count than requested* (e.g. `int16_t` levels with 80000 bins computes as if ~14464 bins had been asked for). The device-side dispatch path performs no such rejection at all.

Scope: `HistogramEven`/`MultiHistogramEven` with integral sample/level types (`HistogramRange` is unaffected).

### How to Reproduce

```cpp
// int16_range_wrap_repro.cu
#include

#include
#include

#include

int main()
{
constexpr int num_levels = 101; // 100 bins of width 655.35 over [-32768, 32767)
const auto d_samples = thrust::device_vector{int16_t{-32768}, int16_t{0}, int16_t{32766}};
// expected bins: 0, 50, 99
thrust::device_vector d_histogram(num_levels - 1, 0);

size_t temp_bytes = 0;
cub::DeviceHistogram::HistogramEven(nullptr, temp_bytes,
thrust::raw_pointer_cast(d_samples.data()), thrust::raw_pointer_cast(d_histogram.data()),
num_levels, int16_t{-32768}, int16_t{32767}, 3);
thrust::device_vector d_temp(temp_bytes);
const cudaError_t err = cub::DeviceHistogram::HistogramEven(thrust::raw_pointer_cast(d_temp.data()), temp_bytes,
thrust::raw_pointer_cast(d_samples.data()), thrust::raw_pointer_cast(d_histogram.data()),
num_levels, int16_t{-32768}, int16_t{32767}, 3);
cudaDeviceSynchronize();
printf("HistogramEven returned: %d (%s)\n", err, cudaGetErrorString(err));

thrust::host_vector h = d_histogram;
for (int k = 0; k < num_levels - 1; ++k)
{
if (h[k] != 0) { printf(" bin %3d: count %d\n", k, h[k]); }
}
printf("expected: bins 0, 50, 99 with count 1 each\n");
return 0;
}
```

From a CCCL checkout:

```bash
nvcc -std=c++17 -arch=sm_89 -Icub -Ithrust -Ilibcudacxx/include int16_range_wrap_repro.cu -o repro
./repro
```

For the int32 flavor, change the sample/level type to `int32_t` with `lower = -1500000000`, `upper = 1500000000` (samples `{-1500000000, 0, 1400000000}`, expected bins `{0, 50, 96}`): the call returns `cudaErrorInvalidValue` and no histogram is computed.

### Expected behavior

A histogram over any valid level range of the sample type should either be computed correctly or, where the documented overflow condition genuinely applies, be rejected with `cudaErrorInvalidValue` — never silently miscounted. For the reproducer: bins 0, 50 and 99 should each hold one sample. Observed output — success is reported and every sample is in bin 0:

```
HistogramEven returned: 0 (no error)
bin 0: count 3
expected: bins 0, 50, 99 with count 1 each
```

For the int32 flavor: the histogram should be computed (bins 0, 50, 96), since the documented rejection condition does not apply. Observed: `cudaErrorInvalidValue`.

### Reproduction link

_No response_

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_

Start by running the supplied int16_t and int32_t reproducers, then inspect cub/device/device_histogram.cuh at the HistogramEven and MultiHistogramEven entry points. Review the existing bin-count overflow coverage introduced by #6908 and trace the integral level-range and bin-count handling. Done means valid ranges produce the expected bins, documented overflow cases are rejected, and invalid bin counts are consistently rejected.

Written by the indexing model from the issue text.