CUDA述語を満たす場合にのみ、変換された結果をコピーする

Question

入力をthrust::device_vectorに変換し、結果が述語を満たす場合は出力ベクトルにコピーします。したがって、結果の数は入力デバイス_ベクトルのサイズよりも小さくなります（出力ベクトルthrust::copy_ifと同様）。私はthrust :: transform_ifでこれを行う方法を見つけていません。以下の例に示すように、現在、私はthrust::transformとthrust::remove_ifでこれを行うことができます。

after adding random number: 
18 4 8 7 11 
after removing values greater than 6: 
4 

私はthrust::transformすることにより、第1、二回のメモリへの結果のコピーを避けたい：出力を与える

#include <thrust/random.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/device_vector.h> 
#include <thrust/transform.h> 
#include <thrust/remove.h> 
#include <iostream> 

__host__ __device__ unsigned int hash(unsigned int a) { 
    a = (a+0x7ed55d16) + (a<<12); 
    a = (a^0xc761c23c)^(a>>19); 
    a = (a+0x165667b1) + (a<<5); 
    a = (a+0xd3a2646c)^(a<<9); 
    a = (a+0xfd7046c5) + (a<<3); 
    a = (a^0xb55a4f09)^(a>>16); 
    return a; 
}; 

struct add_random { 
    __host__ __device__ add_random() {} 
    __device__ int operator()(const int n, const int x) const { 
    thrust::default_random_engine rng(hash(n)); 
    thrust::uniform_int_distribution<int> uniform(0, 11); 
    return uniform(rng)+x; 
    } 
}; 

struct is_greater { 
    __host__ __device__ bool operator()(const int x) { 
    return x > 6 ; 
    } 
}; 

int main(void) { 
    int x[5] = {10, 2, 5, 3, 0}; 
    thrust::device_vector<int> d_x(x, x+5); 

    thrust::transform(
     thrust::counting_iterator<int>(0), 
     thrust::counting_iterator<int>(5), 
     d_x.begin(), 
     d_x.begin(), 
     add_random()); 

    std::cout << "after adding random number:" << std::endl; 
    std::ostream_iterator<int> o(std::cout, " "); 
    thrust::copy(d_x.begin(), d_x.end(), o); 
    std::cout << std::endl; 

    thrust::device_vector<int>::iterator new_end(thrust::remove_if(d_x.begin(), d_x.end(), is_greater())); 

    std::cout << "after removing values greater than 6:" << std::endl; 
    thrust::copy(d_x.begin(), new_end, o); 
    std::cout << std::endl; 

    return 0; 
} 

上記の例ではthrust::remove_ifとなります。 1つの変換関数で上記の出力を得ることは可能ですか？これどうやってするの？私の最大の関心事は計算コストです。したがって、Thrustライブラリを使用しなくても、最適化されたソリューションはすばらしいでしょう。

Answer 1

スラストファンシーイテレータの世界へようこそ。 thrust quick start guideを見ることで、いくつかのイテレータタイプの素早い概要を知ることができます。特に、スラスト変換イテレータは、別の推力アルゴリズムの入力に適用される推力変換操作を置き換えるために頻繁に使用され、2つのアルゴリズムを単一の操作に「融合」します。

ここではあなたのケースに適用される加工された例があります：私たちは同じ2つの入力に適用される変換イテレータを使用して変換操作を交換しました

$ cat t1254.cu 
#include <thrust/random.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/device_vector.h> 
#include <thrust/transform.h> 
#include <thrust/remove.h> 
#include <iostream> 

__host__ __device__ unsigned int hash(unsigned int a) { 
    a = (a+0x7ed55d16) + (a<<12); 
    a = (a^0xc761c23c)^(a>>19); 
    a = (a+0x165667b1) + (a<<5); 
    a = (a+0xd3a2646c)^(a<<9); 
    a = (a+0xfd7046c5) + (a<<3); 
    a = (a^0xb55a4f09)^(a>>16); 
    return a; 
}; 

struct add_random : public thrust::unary_function<thrust::tuple<int, int>, int> { 
    __host__ __device__ int operator()(thrust::tuple<int, int> t) const { 
    int n = thrust::get<0>(t); 
    int x = thrust::get<1>(t); 
    thrust::default_random_engine rng(hash(n)); 
    thrust::uniform_int_distribution<int> uniform(0, 11); 
    return uniform(rng)+x; 
    } 
}; 

struct is_greater { 
    __host__ __device__ bool operator()(const int x) { 
    return x < 6 ; 
    } 
}; 

int main(void) { 
    int x[5] = {10, 2, 5, 3, 0}; 
    thrust::device_vector<int> d_x(x, x+5); 
    thrust::device_vector<int> d_r(5); 
    int rsize = thrust::copy_if(thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator<int>(0), d_x.begin())), add_random()), thrust::make_transform_iterator(thrust::make_zip_iterator(thrust::make_tuple(thrust::counting_iterator<int>(5), d_x.end())), add_random()), d_r.begin(), is_greater())- d_r.begin(); 
    std::cout << "after removing values greater than 6:" << std::endl; 
    thrust::copy_n(d_r.begin(), rsize, std::ostream_iterator<int>(std::cout, " ")); 
    std::cout << std::endl; 

    return 0; 
} 
$ nvcc -o t1254 t1254.cu 
$ ./t1254 
after removing values greater than 6: 
4 
$ 

1. 。変換操作には2つの入力があるため、これらを結合するにはzipイテレーターを使用しています。また、変換ファンクターも入力としてそのタプルを受け入れるために若干修正されています。

2. トランスフォームイテレータを入力として使用するには、remove_ifをcopy_ifに変換します。これには、コピー述部のロジックにわずかな変更が必要です。