CUDA推論sort_by_keyキーがzip_iteratorのカスタム比較述語で処理されたタプルの場合

Question

私はここで多くの質問を似たように見てきましたが、1つのマイナーチェンジがあります。私は複合キーとしてzip_iteratorで値をソートしようとしています。

具体的には、私は以下の機能を持っている：

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

任意のアイデア何かもしれませんこの/私は実際に動作する述語をどのように書くのですか？事前に

おかげで、 ネイサン

Answer 1

コンパレータはタイプconst thrust::tuple<unsigned int, float>&の引数を取ります。あなたが定義されconst tdp_uif_tuple&タイプは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; 

    } 
}; 

は、それは同様にあなたのためにありません願っています:)

http://code.google.com/p/thrust/wiki/QuickStartGuide#zip_iteratorはジップイテレータの詳細を持っています。

必要がありますが、これらのテンプレートの長さをクリーンアップするために探しているなら、あなたはこれを行うことができない。

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が使用され、カスタマイズされた比較演算子である場合のsort_by_keyの使用方法に関する完全な実例です。

#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]); 
    } 

}