[BUG]: [BUG]: cub::DeviceHistogram mis-counts bins and writes out of bounds when the number of bins exceeds what the sample type can represent exactly

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

CUB

### Describe the bug

When the number of bins exceeds the largest integer the *sample type* can represent exactly, `DeviceHistogram` silently produces corrupted counts — and, for floating-point sample types, writes past the end of the output histogram. All calls return `cudaSuccess`. The observable behavior by sample type:

- **`__nv_bfloat16` with more than 256 bins**: neighboring high-index bins are merged — odd bins above index 256 report 0 while their even neighbors double- or triple-count (verified: with 300 bins and exactly one sample per bin, 33 of 300 bins have wrong counts). Additionally, counts destined for the topmost bins can be **written one element past the end of the output histogram** (out-of-bounds `atomicAdd`; verified — one of the 300 counts leaves the histogram entirely).
- **`__half` with more than 2048 bins**: same merging/out-of-bounds behavior, starting at bin index 2048.
- **`int16_t` with more than 32768 bins**: counts for bins above index 32767 are **silently dropped** (verified: samples belonging to bins 32768 and 64768 of a 65535-bin histogram simply vanish). This is a reachable configuration — a full-range `int16_t` histogram legitimately has up to 65535 bins.
- **`float` with more than 2^24 bins**: same failure shape in principle, though such configurations are impractical.
- Unaffected: 8-bit types, `uint16_t`, 32/64-bit integers with realistic bin counts, and `double`.

**`HistogramEven`, `HistogramRange`, and the Multi variants are all affected.** Notably this corrupts `HistogramRange` even though its binary-search binning computes every bin index exactly — the corruption happens after binning, when per-block partial results are folded into the output histogram, where each bin index makes a round trip through the sample type (`cub/agent/agent_histogram.cuh`, `StoreOutput`) and is altered by rounding (floating-point types) or overflow (`int16_t`).

The same round-trip also **fails to compile** for `__nv_bfloat16` on CTK 12.0 ("more than one constructor applies to convert from `int`"), so on that toolkit the bf16 corruption is a build break instead.

Related but distinct: #10975 covers the bin *computation* going wrong when the level range exceeds the sample type's capacity; this issue is about the bin *counts* being misplaced after a correct computation. Found while adding `__nv_bfloat16` test coverage for #10940, but independent of that fix.

### How to Reproduce

`__nv_bfloat16`, isolating the failure via `HistogramRange` (exact search binning, 300 bins, one sample placed in each bin):

```cpp
// bf16_roundtrip_repro.cu
#include

#include
#include

#include

#include
#include

int main()
{
constexpr int num_bins = 300;
constexpr int num_levels = num_bins + 1;

// 301 consecutive bf16 values starting at 1.0: exactly representable, strictly increasing
std::vector<__nv_bfloat16> h_levels(num_levels);
for (int k = 0; k < num_levels; ++k)
{
h_levels[k] = __ushort_as_bfloat16(static_cast(0x3F80 + k));
}
// one sample per bin, exactly on the bin's lower edge -> sample k belongs to bin k
std::vector<__nv_bfloat16> h_samples(h_levels.begin(), h_levels.begin() + num_bins);

thrust::device_vector<__nv_bfloat16> d_levels(h_levels);
thrust::device_vector<__nv_bfloat16> d_samples(h_samples);
thrust::device_vector d_histogram(num_bins, 0);

size_t temp_bytes = 0;
cub::DeviceHistogram::HistogramRange(
nullptr, temp_bytes,
thrust::raw_pointer_cast(d_samples.data()), thrust::raw_pointer_cast(d_histogram.data()),
num_levels, thrust::raw_pointer_cast(d_levels.data()), num_bins);
thrust::device_vector d_temp(temp_bytes);
cub::DeviceHistogram::HistogramRange(
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, thrust::raw_pointer_cast(d_levels.data()), num_bins);
cudaDeviceSynchronize();

thrust::host_vector h_histogram = d_histogram;
int wrong = 0, total = 0;
for (int k = 0; k < num_bins; ++k)
{
total += h_histogram[k];
if (h_histogram[k] != 1)
{
printf(" bin %3d: count %d (expected 1)\n", k, h_histogram[k]);
++wrong;
}
}
printf("bins with wrong counts: %d of %d (total counted samples: %d of %d)\n", wrong, num_bins, total, num_bins);
printf(wrong == 0 ? "PASS\n" : "FAIL\n");
return wrong == 0 ? 0 : 1;
}
```

From a CCCL checkout:

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

The `int16_t` variant (silent drops instead of merges): use 65536 levels `-32768..32767` (65535 bins) and samples `{-32768, 0, 32000}`, which belong in bins `{0, 32768, 64768}`. Only bin 0 receives its count; the other two samples vanish.

### Expected behavior

Every bin should contain exactly one sample. Observed output — odd bins above index 256 are emptied, their neighbors double/triple-count, and one sample is written out of bounds (total counted is 299 of 300):

```
bin 256: count 2 (expected 1)
bin 257: count 0 (expected 1)
bin 259: count 0 (expected 1)
bin 260: count 3 (expected 1)
bin 261: count 0 (expected 1)
bin 263: count 0 (expected 1)
bin 264: count 3 (expected 1)
...
bins with wrong counts: 33 of 300 (total counted samples: 299 of 300)
FAIL
```

The missing 300th count is written one element past the end of the output histogram.

### Reproduction link

_No response_

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_

Start with cub/agent/agent_histogram.cuh, especially StoreOutput, and trace how DeviceHistogram::HistogramRange folds per-block results into the output histogram. Run the provided __nv_bfloat16 reproduction, then the int16_t variant, and verify that all bins retain their counts without out-of-bounds writes across the affected HistogramEven, HistogramRange, and Multi variants.

Written by the indexing model from the issue text.