在设备阵列上按键减少答案

【问题标题】：Reduce by key on device array在设备阵列上按键减少
【发布时间】：2015-08-27 09:05:17
【问题描述】：

我正在使用 reduce_by_key 来查找 int2 类型数组中的元素数，该数组具有相同的第一个值。

例如数组：所以不行。以 1 为第一个元素的元素是 3，以 2 为首元素的元素是 2。

代码：

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;
}

上面的代码显示 Segmentation Fault 。我使用 gdb 运行了上面的代码，它报告了这个位置的段错误。

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=...)

在 /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

如何在设备阵列上使用 reduce_by_key ？

【问题讨论】：

标签： cuda parallel-processing thrust

【解决方案1】：

Thrust 将普通指针解释为指向主机上的数据：

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

因此，thrust 将为上述算法调用主机路径，当它试图取消引用主机代码中的这些指针时，它会出现段错误。这在thrust getting started guide:

您可能想知道当“原始”指针用作 Thrust 函数的参数时会发生什么。与 STL 一样，Thrust 允许这种用法，它将分派算法的主机路径。如果有问题的指针实际上是指向设备内存的指针，那么您需要在调用函数之前用推力::device_ptr 包装它。

Thrust 有多种机制（例如device_ptr、device_vector 和执行策略）来向算法识别数据是设备驻留的并且应该使用设备路径。

对现有代码的最简单修改可能是使用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);

上述问题与another issue you asked about类似。

【讨论】：