CUDA Thrust sort_by_key when the key is a tuple dealt with by zip_iterator's with custom comparison predicate

Question

I've looked through a lot of questions here for something similar and there are quite a few, albeit with one minor change. I'm trying to sort values with a zip_iterator as a compound key.

Specifically, I have the following function:

void thrustSort(
    unsigned int * primaryKey,
    float * secondaryKey,
    unsigned int * values,
    unsigned int numberOfPoints)
{
    thrust::device_ptr dev_ptr_pkey = thrust::device_pointer_cast(primaryKey);
    thrust::device_ptr dev_ptr_skey = thrust::device_pointer_cast(secondaryKey); 
    thrust::device_ptr dev_ptr_values = thrust::device_pointer_cast(values);

    thrust::tuple,thrust::device_ptr> keytup_begin =
        thrust::make_tuple,thrust::device_ptr>(dev_ptr_pkey, dev_ptr_skey);

    thrust::zip_iterator, thrust::device_ptr > > first =
        thrust::make_zip_iterator, thrust::device_ptr > >(keytup_begin);

    thrust::sort_by_key(first, first + numberOfPoints, dev_ptr_values, ZipComparator());    
}

and this custom predicate:

typedef thrust::device_ptr<unsigned int> tdp_uint ;
typedef thrust::device_ptr<float> tdp_float ;
typedef thrust::tuple<tdp_uint, tdp_float> tdp_uif_tuple ;

struct ZipComparator
{
    __host__ __device__
    inline bool operator() (const tdp_uif_tuple &a, const tdp_uif_tuple &b)
    {
        if(a.head < b.head) return true;
        if(a.head == b.head) return a.tail < b.tail;
        return false;

    }
};

The errors I'm getting are:

Error   1   error : no instance of constructor "thrust::device_ptr::device_ptr [with T=unsigned int]" matches the argument list  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include\thrust\detail\tuple.inl 309 1   ---
Error   2   error : no instance of constructor "thrust::device_ptr::device_ptr [with T=float]" matches the argument list C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include\thrust\detail\tuple.inl 401 1   ---

Any ideas what might cause this / how do I write a predicate that indeed works?

Thanks in Advance, Nathan

Answer 1

This is a fully worked example on how using sort_by_key when the key is a tuple dealt with by zip_iterator's and a customized comparison operator.

#include <thrust/device_vector.h>
#include <thrust/sort.h>

#include "Utilities.cuh"

// --- Defining tuple type
typedef thrust::tuple<int, int> Tuple;

/**************************/
/* TUPLE ORDERING FUNCTOR */
/**************************/
struct TupleComp
{
    __host__ __device__ bool operator()(const Tuple& t1, const Tuple& t2)
    {
        if (t1.get<0>() < t2.get<0>())
            return true;
        if (t1.get<0>() > t2.get<0>())
            return false;
        return t1.get<1>() < t2.get<1>();
    }
};

/********/
/* MAIN */
/********/
int main()
{
    const int N = 8;

    // --- Keys and values on the host: allocation and definition
    int h_keys1[N]      = { 1, 3, 3, 3, 2, 3, 2, 1 };                                         
    int h_keys2[N]      = { 1, 5, 3, 8, 2, 8, 1, 1 };                                         
    float h_values[N]   = { 0.3, 5.1, 3.2, -0.08, 2.1, 5.2, 1.1, 0.01};

    printf("\n\n");
    printf("Original\n");
    for (int i = 0; i < N; i++) {
        printf("%i %i %f\n", h_keys1[i], h_keys2[i], h_values[i]);
    }

    // --- Keys and values on the device: allocation
    int *d_keys1;       gpuErrchk(cudaMalloc(&d_keys1, N * sizeof(int)));
    int *d_keys2;       gpuErrchk(cudaMalloc(&d_keys2, N * sizeof(int)));
    float *d_values;    gpuErrchk(cudaMalloc(&d_values, N * sizeof(float)));

    // --- Keys and values: host -> device
    gpuErrchk(cudaMemcpy(d_keys1, h_keys1, N * sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_keys2, h_keys2, N * sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_values, h_values, N * sizeof(float), cudaMemcpyHostToDevice));

    // --- From raw pointers to device_ptr
    thrust::device_ptr<int> dev_ptr_keys1 = thrust::device_pointer_cast(d_keys1);
    thrust::device_ptr<int> dev_ptr_keys2 = thrust::device_pointer_cast(d_keys2);
    thrust::device_ptr<float> dev_ptr_values = thrust::device_pointer_cast(d_values);

    // --- Declare outputs
    thrust::device_vector<float> d_values_output(N);
    thrust::device_vector<Tuple> d_keys_output(N);

    auto begin_keys = thrust::make_zip_iterator(thrust::make_tuple(dev_ptr_keys1, dev_ptr_keys2));
    auto end_keys = thrust::make_zip_iterator(thrust::make_tuple(dev_ptr_keys1 + N, dev_ptr_keys2 + N));

    thrust::sort_by_key(begin_keys, end_keys, dev_ptr_values, TupleComp());

    int *h_keys1_output = (int *)malloc(N * sizeof(int));
    int *h_keys2_output = (int *)malloc(N * sizeof(int));
    float *h_values_output = (float *)malloc(N * sizeof(float));

    gpuErrchk(cudaMemcpy(h_keys1_output, d_keys1, N * sizeof(int), cudaMemcpyDeviceToHost));
    gpuErrchk(cudaMemcpy(h_keys2_output, d_keys2, N * sizeof(int), cudaMemcpyDeviceToHost));
    gpuErrchk(cudaMemcpy(h_values_output, d_values, N * sizeof(float), cudaMemcpyDeviceToHost));

    printf("\n\n");
    printf("Ordered\n");
    for (int i = 0; i < N; i++) {
        printf("%i %i %f\n", h_keys1_output[i], h_keys2_output[i], h_values_output[i]);
    }

}

Answer 2

The comparator takes arguments of type const thrust::tuple<unsigned int, float>&. The const tdp_uif_tuple& type you defined expands to const thrust::tuple<thrust::device_ptr<unsigned int>, thrust:device_ptr<float> >&

The code below compiles for me:

struct ZipComparator
{
    __host__ __device__
    inline bool operator() (const thrust::tuple<unsigned int, float> &a, const thrust::tuple<unsigned int, float> &b)
    {
        if(a.head < b.head) return true;
        if(a.head == b.head) return a.tail < b.tail;
        return false;

    }
};

Hope it does for you as well :)

http://code.google.com/p/thrust/wiki/QuickStartGuide#zip_iterator has more details on the zip iterator.

Not required, but if you're looking to clean up the length of those templates, you can do this:

void thrustSort(
    unsigned int * primaryKey,
    float * secondaryKey,
    unsigned int * values,
    unsigned int numberOfPoints)
{
    tdp_uint dev_ptr_pkey(primaryKey);
    tdp_float dev_ptr_skey(secondaryKey);   
    tdp_uint dev_ptr_values(values);

    thrust::tuple<tdp_uint, tdp_float> keytup_begin = thrust::make_tuple(dev_ptr_pkey, dev_ptr_skey);

    thrust::zip_iterator<thrust::tuple<tdp_uint, tdp_float> > first =
    thrust::make_zip_iterator(keytup_begin);

    thrust::sort_by_key(first, first + numberOfPoints, dev_ptr_values, ZipComparator());    
}

A lot of the template arguments can be inferred from the arguments.