Fix select if for mixed types

[gevtushenko]

During three-way partitioning development I've noticed that there's a bug in the select if implementation. It used output type to load data from input iterator. Recently, @elstehle brought the very same issue. Here's a code sample that fails to compile. This fact limits the applicability of the partitioning/selection facilities:

#include <iostream>

#include <thrust/count.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/transform_output_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/device_vector.h>

#include <cub/device/device_select.cuh>

struct pair_to_col {
    __host__ __device__ int operator()(const thrust::tuple<int, double> &in) {
        return thrust::get<0>(in);
    }
};

struct select {
    __host__ __device__ bool operator()(const thrust::tuple<int, double> &in) {
        return static_cast<double>(thrust::get<0>(in)) > thrust::get<1>(in);
    }
};

int main() {
    const std::size_t n = 1024;
    thrust::device_vector<int> col_a(n, 42);
    thrust::device_vector<double> col_b(n, 4.2);

    thrust::device_vector<int> result(n, 42);

    auto in = thrust::make_zip_iterator(col_a.begin(), col_b.begin());
    auto out = thrust::make_transform_output_iterator(result.begin(), pair_to_col{});

    std::size_t tmp_storage_size{};
    cub::DeviceSelect::If(
            nullptr, tmp_storage_size, in, out, thrust::make_discard_iterator(), n, select{});

    thrust::device_vector<char> tmp_storage(tmp_storage_size);
    cub::DeviceSelect::If(
            thrust::raw_pointer_cast(tmp_storage.data()), tmp_storage_size, in, out, thrust::make_discard_iterator(), n, select{});

    if (thrust::count(result.begin(), result.end(), 42) == n) {
        std::cout << "OK" << std::endl;
    } else {
        std::cout << "FAIL" << std::endl;
    }

    return 0;
}

This PR addresses this issue, allowing different value types for input and output iterators. It's worth mentioning that the implementation used to have correct types up to the following commit which doesn't contain any mentions of the value types in its description. So my guess whould be that it was a typo rather than an intent.

I've also noticed that the radix sort test fails to compile on windows. I've fixed a few conversion issues here.

[alliepiper]

@senior-zero Sounds like this may also fix #406, can you confirm?

[elstehle]

Thanks for jumping onto this so quickly, @senior-zero! LGTM. This will definitely help my use case. Given that the change also follows the behaviour we have in the prefix scan, I think it makes sense too. I'm curious if, originally, making InputIt::value_type the internal_type was a conscious decision and if there's rationale behind.

[elstehle]

Good catch 👍 I probably would have missed adjusting the tuning policies.

[gevtushenko]

@senior-zero Sounds like this may also fix #406, can you confirm?

Yes, I confirm. After the fix, the code from #406 gets compiled.

[alliepiper]

LGTM -- I'll start tests in a moment.

[alliepiper]

gpuCI: NVIDIA/thrust#1633
DVS CL: 31093334