Question

我在座谈了许多类似问题,但有一些小小小改动。我试图将价值与一种焦炭_化器混为一谈,将其作为复合钥匙。

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

以及这一习俗:

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;

    }
};

错误是:

Error   1   error : no instance of constructor "thrust::device_ptr::device_ptr [with T=unsigned int]" matches the argument list  C:Program FilesNVIDIA GPU Computing ToolkitCUDAv4.0include	hrustdetail	uple.inl 309 1   ---
Error   2   error : no instance of constructor "thrust::device_ptr::device_ptr [with T=float]" matches the argument list C:Program FilesNVIDIA GPU Computing ToolkitCUDAv4.0include	hrustdetail	uple.inl 401 1   ---

任何可能产生这种结果的想法,我如何写出确实行之有效的前提?

Thanks in Advance, Nathan

Answer 1

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> >&

下面的法典对我进行了汇编:

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;

    }
};

希望也是你们的:

关于燃烧炉的更多细节。

如果你重新寻找这些模板的长度,你可以这样做:

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

可以从这些论点推断出许多模板论点。

Answer 2

这是关于如何使用<条码>的完全有效的实例。当钥匙是tuple,通过zip_iterator<>/code>处理,以及定制的比较操作者。



#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("

");
    printf("Original
");
    for (int i = 0; i < N; i++) {
        printf("%i %i %f
", 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("

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

}

友情链接