[BUG]: cub::DeviceReduce::Sum is not deterministic when reducing a sub range (AA[BB]) compared to the extracted range alone ([BB])

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

Runtime Error

### Component

CUB

### Describe the bug

I'm seeing different results when running cub::DeviceReduce::Sum on a sub range (AA[BB]) compared to the extracted range alone ([BB])

More detailed:

Input is a vector of two parts, A and B, both of the same size. We only care about B
Goal is to highlight a difference in results when using cub::DeviceReduce::Sum on B directly vs B as part of the two part vector
Visually:
Two_part_vector = [A, A, B, B]
One_part_vector = [B, B]
Sum(Two_part_vector + 2) == Sum(One_part_vector)

But in practice the sums have a 1 bit difference.

### How to Reproduce

Results were found on both L40, nvcc 12.8 and -arch=native (without, we get the same results) and RTX5090 with nvcc 13.0 and *not* using -arch=native (-arch=native make it work)

Code:
```c++
#include
#include
#include
#include
#include
#include
#include
#include
#include
#include

// Creates a vector from a binary file
template
std::vector read_vector_from_file(const std::string& filename)
{
std::ifstream in_file(filename, std::ios::binary | std::ios::in);
if (!in_file.is_open()) {
std::cerr << "Failed to open file for reading: " << filename << std::endl;
exit(-1);
return std::vector();
}

std::size_t size;
in_file.read(reinterpret_cast(&size), sizeof(size));

std::vector vec(size);
in_file.read(reinterpret_cast(vec.data()), size * sizeof(T));
in_file.close();

return vec;
}

// Converts a double to a uint64_t (to have a bit wise comparable output)
__host__ __device__ uint64_t to_long_long(double input)
{
uint64_t ll;
memcpy(&ll, &input, sizeof(input));
return ll;
}

// Input is a vector of two parts, A and B, both of the same size. We only care about B
// Goal is to create a non determistic results when using cub::DeviceReduce::Sum on B directly vs B as part of the two part vector
// Visually:
// Two_part_vector = [A, A, B, B]
// One_part_vector = [B, B]
// Sum(Two_part_vector + 2) == Sum(One_part_vector)
int main()
{
// Load the values vector
thrust::host_vector vec = read_vector_from_file("bound_value_batch_all.bin");
std::cout << "Input size " << vec.size() << std::endl;
int semi_size = vec.size() / 2;

// Allocate the vector that will contain just B
thrust::device_vector vector_exact_size(semi_size);
// Copy the B part of the vector
thrust::copy(vec.begin() + semi_size, vec.end(), vector_exact_size.begin());

// Allocate the vector that will contain A and B
double* vector_twice_the_size;
cudaMalloc(&vector_twice_the_size, vec.size() * sizeof(double));
cudaMemcpy(vector_twice_the_size, vec.data(), vec.size() * sizeof(double), cudaMemcpyHostToDevice);

// Test exact equality between both vectors, B and the B part of the two part vector
const bool equal = thrust::equal(thrust::device,
vector_twice_the_size + semi_size,
vector_twice_the_size + vec.size(),
vector_exact_size.begin());
if (equal) {
std::cout << "Both vectors (B and B part of the two part vector) are equal" << std::endl;
} else {
std::cout << "Both vectors (B and B part of the two part vector) are different" << std::endl;
}

// Prepare the code to run the two cub::DeviceReduce::Sum
double* out1;
double* out2;
cudaMalloc(&out1, sizeof(double));
cudaMalloc(&out2, sizeof(double));

// Allocate temporary memory for cub::DeviceReduce::Sum
uint8_t* tmp_1 = nullptr;
uint8_t* tmp_2 = nullptr;
size_t size_of_tmp_1 = 0;
size_t size_of_tmp_2 = 0;

// FIXME: when .begin() is used, the sum is correct. If a raw pointer is used, the sum is wrong.
auto vector_exact_size_start = thrust::raw_pointer_cast(vector_exact_size.data());
//auto vector_exact_size_start = vector_exact_size.begin();

cub::DeviceReduce::Sum(tmp_1, size_of_tmp_1, vector_exact_size_start, out1, semi_size);
cub::DeviceReduce::Sum(tmp_2, size_of_tmp_2, vector_twice_the_size + semi_size, out2, semi_size);

cudaMalloc(&tmp_1, size_of_tmp_1);
cudaMalloc(&tmp_2, size_of_tmp_2);

// Run the two cub::DeviceReduce::Sum
cub::DeviceReduce::Sum(tmp_1, size_of_tmp_1, vector_exact_size_start, out1, semi_size);
cub::DeviceReduce::Sum(tmp_2, size_of_tmp_2, vector_twice_the_size + semi_size, out2, semi_size);

// Return the results
double out1_host;
double out2_host;
cudaMemcpy(&out1_host, out1, sizeof(double), cudaMemcpyDeviceToHost);
cudaMemcpy(&out2_host, out2, sizeof(double), cudaMemcpyDeviceToHost);

// Print the results (using uint64_t to avoid precision issues)
std::cout << "out1_host=" << to_long_long(out1_host) << std::endl;
std::cout << "out2_host=" << to_long_long(out2_host) << std::endl;
if (out1_host == out2_host) {
std::cout << "The two results are equal" << std::endl;
} else {
std::cout << "The two results are different" << std::endl;
}

cudaFree(vector_twice_the_size);
cudaFree(out1);
cudaFree(out2);
cudaFree(tmp_1);
cudaFree(tmp_2);
return 0;
}
```

Binary vector used in the example is shared attached. It's rare and hard to find a vector which triggers the same issue.

[bound_value_batch_all.zip](https://github.com/user-attachments/files/21733696/bound_value_batch_all.zip)

### Expected behavior

Results of cub::DeviceReduce::Sum should be deterministic in all cases, no matter where the memory is and if it's a .data or .begin

### Reproduction link

_No response_

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_