推力ファクタ内で推力アルゴリズムを呼び出す

Question

thrust :: transform_reduceの引数であるファンクタ内で、thrust :: reduceを使用しています。状況は、入れ子状の推力アルゴリズムのように見えます。コンパイルは成功しますが、それはエラーで実行されます。

terminate called after throwing an instance of 'thrust::system::system_error' 
    what(): cudaEventSynchronize in future::wait: an illegal memory access was encountered 
Aborted (core dumped) 

以下のようにコードは次のとおりです。

#include <thrust/inner_product.h> 
#include <thrust/functional.h> 
#include <thrust/device_vector.h> 

#include <iostream> 
#include <cmath> 
#include <boost/concept_check.hpp> 


struct aFuntor : public thrust::unary_function<int, int> 
{ 
    aFuntor(int* av__, int* bv__, const int& N__) : av_(av__), bv_(bv__), N_(N__) {}; 

    __host__ __device__ 
    int operator()(const int& idx) 
    { 

    thrust::device_ptr<int> av_dpt = thrust::device_pointer_cast(av_); 

    int res = thrust::reduce(av_dpt, av_dpt+N_); 

     return res; 
    } 

    int* av_; 
    int* bv_; 
    int N_; 
}; 


int main(void) 
{ 
     int N = 5; 
     std::vector<int> av = {0,1,3,5}; 
     std::vector<int> bv = {0,10,20,30}; 
     thrust::device_vector<int> av_d(N); 
     thrust::device_vector<int> bv_d(N); 
     av_d = av; bv_d = bv; 

     // initial value of the reduction 
     int init=0; 

     // binary operations 
     thrust::plus<int>  bin_op; 

     int res = 
     thrust::transform_reduce(thrust::counting_iterator<int>(0), 
           thrust::counting_iterator<int>(N-1), 
        aFuntor(thrust::raw_pointer_cast(av_d.data()), 
         thrust::raw_pointer_cast(bv_d.data()), 
         N), 
       init, 
       bin_op);  

     std::cout << "result is: " << res << std::endl; 
     return 0; 
} 

は、ネストされた、この種の構造をサポートして推力のでしょうか？または私のアルゴリズムを再設計する必要がある以外の方法はありませんか？並列化が困難なアルゴリズムが存在しますか？

ありがとうございます！

Answer 1

推力はnested algorithm usageを許容します。しかし、スラストがデバイスコードから推力アルゴリズムを起動するときにスラストがデバイスパスのみを選択することを確認する必要があります。この場合、これは起こっていません。そのまま私はあなたのコードをコンパイルするときに、少なくとも私のシステム（Ubuntuの14.04）上で、私はそれの表示を得る：

t113.cu(20) (col. 9): warning: calling a __host__ function("thrust::reduce< ::thrust::device_ptr<int> > ") from a __host__ __device__ function("aFuntor::operator()") is not allowed 

だから、ここに望まれるものを明確ではありません。代わりに、推力の実行ポリシーがthrust::deviceのデバイスパス（デバイスコードでは、デバイスポインタを渡しているため、ファンクタの定義に本質的に暗黙的に含まれています）を使用することができます。私は、以下の変更を加えた場合、あなたのコードは私のために、エラーなしでコンパイルして実行します：

$ cat t113.cu 
#include <thrust/inner_product.h> 
#include <thrust/functional.h> 
#include <thrust/device_vector.h> 

#include <iostream> 
#include <cmath> 
#include <thrust/execution_policy.h> 
//#include <boost/concept_check.hpp> 


struct aFuntor : public thrust::unary_function<int, int> 
{ 
    aFuntor(int* av__, int* bv__, const int& N__) : av_(av__), bv_(bv__), N_(N__) {}; 

    __host__ __device__ 
    int operator()(const int& idx) 
    { 

    thrust::device_ptr<int> av_dpt = thrust::device_pointer_cast(av_); 

    int res = thrust::reduce(thrust::device, av_dpt, av_dpt+N_); 

     return res; 
    } 

    int* av_; 
    int* bv_; 
    int N_; 
}; 


int main(void) 
{ 
     int N = 5; 
     std::vector<int> av = {0,1,3,5}; 
     std::vector<int> bv = {0,10,20,30}; 
     thrust::device_vector<int> av_d(N); 
     thrust::device_vector<int> bv_d(N); 
     av_d = av; bv_d = bv; 

     // initial value of the reduction 
     int init=0; 

     // binary operations 
     thrust::plus<int>  bin_op; 

     int res = 
     thrust::transform_reduce(thrust::counting_iterator<int>(0), 
           thrust::counting_iterator<int>(N-1), 
        aFuntor(thrust::raw_pointer_cast(av_d.data()), 
         thrust::raw_pointer_cast(bv_d.data()), 
         N), 
       init, 
       bin_op); 

     std::cout << "result is: " << res << std::endl; 
     return 0; 
} 
$ nvcc -std=c++11 -arch=sm_61 -o t113 t113.cu 
$ ./t113 
result is: 36 
$ 

私は実際にコードから自分の意図を解析するために試したことがないので、私はこれが正しいことを確認するために言うことができません答えは、それはあなたが求めている質問ではないようです。 （後で答えが正しいと思われます。あなたのファンクタは、すべての要素に対して9の値を生成しています。そして、4要素9で9を減らしています。

すべてのことを言っても、なぜ推力が元のケースでホストパスを選択しているのかは完全にはわかりません。あなたが好きなら、そのためにthrust issueを提出することができます。しかし、スラストディスパッチシステムを十分に注意深く考えなかったことは全く可能です。ホストコードアルゴリズムのディスパッチ（transform_reduce）は、たとえば、ホストコンテナまたはデバイスコンテナを使用しているかどうかがわからない場合があるため、やや混乱するかもしれません。