[NVBug 3298282] Device-side CDP launch of `thrust::upper_bound` returns incorrect result.

Reported via NVBug 3298282. Reproduced locally. Appears to be a regression (see comments in test case). Only happens when launching from the device, suggesting that the issue is in the sequential `binary_search` implementation.

```
// simple Thrust failure reproducer by Alex Long (along@lanl.gov)
// passes with CUDA 11.0,2 upper_bound returns 3.0
// fails with CUDA 11.2, upper_bound returns 0.0

#include
#include
#include
#include
#include
#include

constexpr double search_value = 2.0;
constexpr int n_blocks = 1;
constexpr int n_threads_per_block =1;

__global__ void do_lower_bound(const double search_value, double *bounds, double * upper_bound) {

auto upper_bound_itr = thrust::lower_bound(thrust::device, &bounds[0],
&bounds[0]+6, search_value);
printf("(in kernel) upper_bound is: %f\n", *upper_bound_itr);
upper_bound[0] = *upper_bound_itr;
}

int main(void)
{
// set up bounds and copy to device
std::array bounds = {0.0, 0.1, 0.5, 1.0, 3.0, 5.0};
double *d_bounds;
cudaError_t err = cudaMalloc((void **)&d_bounds, sizeof(double)*bounds.size());
if (err)
std::cout<<"CUDA error in allocating bounds"< upper_bound{0.0};
double *d_upper_bound;
err = cudaMalloc((void **)&d_upper_bound, sizeof(double));
if (err)
std::cout<<"CUDA error in allocating one double"<>>(search_value, d_bounds, d_upper_bound);

auto h_iter = thrust::lower_bound(thrust::device, d_bounds, d_bounds + 6, search_value);
double h_result{};
cudaMemcpy(&h_result, h_iter, sizeof(double), cudaMemcpyDeviceToHost);

cudaDeviceSynchronize();

// copy upper bound value back to host
err = cudaMemcpy(upper_bound.data(), d_upper_bound, sizeof(double), cudaMemcpyDeviceToHost);

auto cpu_upper_bound = *(std::upper_bound(bounds.begin(), bounds.end(), search_value));
std::cout<<"Upper bound should be: "<