[BUG]: cub::DeviceHistogram silently drops all negative int8_t samples (Even and Range)

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

For `int8_t` samples (and `char` on platforms where it is signed), `DeviceHistogram` never counts negative samples: they vanish from the histogram as if they had been outside the level range, and the call returns `cudaSuccess`. Non-negative samples are counted correctly, so the result looks plausible — the total count is just silently too low (verified: 4 samples, one per bin, over levels [-60, 64) → the two negative samples are missing, the two non-negative ones are binned correctly).

**`HistogramEven`, `HistogramRange`, and the Multi variants are all affected**, for any level range that includes negative values. `uint8_t` samples are unaffected.

### How to Reproduce

```cpp
// int8_negative_samples_repro.cu
#include

#include
#include

#include

int main()
{
constexpr int num_levels = 5; // 4 bins of width 31 over [-60, 64)
const auto d_samples = thrust::device_vector{int8_t{-60}, int8_t{-1}, int8_t{10}, int8_t{63}};
// expected bins: 0, 1, 2, 3
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, int8_t{-60}, int8_t{64}, 4);
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, int8_t{-60}, int8_t{64}, 4);
cudaDeviceSynchronize();
printf("HistogramEven returned: %d (%s)\n", err, cudaGetErrorString(err));

thrust::host_vector h = d_histogram;
int total = 0;
for (int k = 0; k < num_levels - 1; ++k)
{
total += h[k];
printf(" bin %d: count %d\n", k, h[k]);
}
printf("total: %d of 4 (expected 1 per bin)\n", total);
return 0;
}
```

From a CCCL checkout:

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

The level range [-60, 64) is chosen so its span (124) fits `int8_t` — this keeps the reproducer independent of #10975 (whose range wrap additionally misbins the *surviving* samples for wider spans, e.g. the full [-128, 127) range).

### Expected behavior

Each of the four samples should land in its own bin. Observed output — both negative samples are silently dropped while the non-negative ones are binned correctly:

```
HistogramEven returned: 0 (no error)
bin 0: count 0
bin 1: count 0
bin 2: count 1
bin 3: count 1
total: 2 of 4 (expected 1 per bin)
```

### Reproduction link

_No response_

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_

Start with cub/device/device_histogram.cuh and compile the supplied int8_negative_samples_repro.cu using the shown nvcc command. Check HistogramEven, HistogramRange, and their Multi variants with negative int8_t samples across the stated level range. Done means all four samples are counted in bins 0 through 3 and the call still returns cudaSuccess.

Written by the indexing model from the issue text.