Reduce by key on device array

Question

I am using reduce_by_key to find the number of elements in an array of type int2 which has same first values .

For example Array: <1,2> <1,3> <1,4> <2,5> <2,7> so no. elements with 1 as first element are 3 and with 2 are 2.

CODE:

struct compare_int2 : public thrust::binary_function<int2, int2, bool> {
__host__ __device__ bool operator()(const int2 &a,const int2 &b) const{
return (a.x == b.x);}
};

compare_int2 cmp;
int main()
{
        int n,i;
        scanf("%d",&n);

        int2 *h_list = (int2 *) malloc(sizeof(int2)*n);
        int *h_ones = (int *) malloc(sizeof(int)*n);

        int2 *d_list,*C;
        int  *d_ones,*D;
        cudaMalloc((void **)&d_list,sizeof(int2)*n);
        cudaMalloc((void **)&d_ones,sizeof(int)*n);
        cudaMalloc((void **)&C,sizeof(int2)*n);
        cudaMalloc((void **)&D,sizeof(int)*n);

        for(i=0;i<n;i++)
        {
                int2 p;
                printf("Value ? ");
                scanf("%d %d",&p.x,&p.y);
                h_list[i] = p;
                h_ones[i] = 1;
        }

        cudaMemcpy(d_list,h_list,sizeof(int2)*n,cudaMemcpyHostToDevice);
        cudaMemcpy(d_ones,h_ones,sizeof(int)*n,cudaMemcpyHostToDevice);
        thrust::reduce_by_key(d_list, d_list+n, d_ones, C, D,cmp);
        return 0;
}

The above code is showing Segmentation Fault . I ran the above code using gdb and it reported the segfault at this location.

thrust::system::detail::internal::scalar::reduce_by_key > (keys_first=0x1304740000,keys_last=0x1304740010,values_first=0x1304740200,keys_output=0x1304740400, values_output=0x1304740600,binary_pred=...,binary_op=...)

at /usr/local/cuda-6.5/bin/../targets/x86_64-linux/include/thrust/system/detail/internal/scalar/reduce_by_key.h:61 61
InputKeyType temp_key = *keys_first

How to use reduce_by_key on device arrays ?

Answer 1

Thrust interprets ordinary pointers as pointing to data on the host:

    thrust::reduce_by_key(d_list, d_list+n, d_ones, C, D,cmp);

Therefore thrust will call the host path for the above algorithm, and it will seg fault when it attempts to dereference those pointers in host code. This is covered in the thrust getting started guide:

You may wonder what happens when a "raw" pointer is used as an argument to a Thrust function. Like the STL, Thrust permits this usage and it will dispatch the host path of the algorithm. If the pointer in question is in fact a pointer to device memory then you'll need to wrap it with thrust::device_ptr before calling the function.

Thrust has a variety of mechanisms (e.g. device_ptr, device_vector, and execution policy) to identify to the algorithm that the data is device-resident and the device path should be used.

The simplest modification for your existing code might be to use device_ptr:

#include <thrust/device_ptr.h>
...
thrust::device_ptr<int2> dlistptr(d_list);
thrust::device_ptr<int>  donesptr(d_ones);
thrust::device_ptr<int2> Cptr(C);
thrust::device_ptr<int>  Dptr(D);
thrust::reduce_by_key(dlistptr, dlistptr+n, donesptr, Cptr, Dptr,cmp);

The issue described above is similar to another issue you asked about.