Aufruf Schubalgorithmen in einem Schub Funktor

Question

Ich benutze Schub :: reduzieren in einem Funktor, die ein Argument in Schub ist :: transform_reduce. Die Situation sieht wie ein verschachtelter Schubalgorithmus aus. Die Zusammenstellung gelingt aber es läuft mit Fehler:

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) 

Der Code wird wie folgt:

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

hat Schub unterstützt diese Art von verschachtelter Struktur? oder es gibt keinen anderen Weg, als meinen Algorithmus neu zu gestalten? AFAIK gibt es Algorithmen, die schwierig sind, Parallelität aufzudecken?

Vielen Dank im Voraus!

Answer 1

Schub ermöglicht nested algorithm usage. Es muss jedoch sichergestellt werden, dass der Schub nur den Gerätepfad auswählt, wenn Schubalgorithmen aus dem Gerätecode gestartet werden, und in Ihrem Fall geschieht dies nicht. Zumindest auf meinem System (Ubuntu 14.04), wenn ich Ihren Code zu kompilieren, wie sie ist, erhalte ich einen Hinweis auf, dass:

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 

das so ist eindeutig nicht das, was hier gesucht. Stattdessen können wir Schub erzwingen, um den Gerätepfad (im Gerätecode - der im Wesentlichen in Ihrer funktor-Definition implizit ist, da Sie Gerätezeiger übergeben) mit einer Schubausführungsrichtlinie von thrust::device zu verwenden. Wenn ich die folgenden Änderungen vornehmen, Ihren Code kompiliert und läuft ohne Fehler für mich:

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

Ich habe nicht wirklich versucht, Ihre Absicht aus dem Code zu analysieren, so dass ich kann nicht sicher sagen, dass dies die richtige ist antworte, aber das scheint nicht die Frage zu sein, die du stellst. (Später: Die Antwort scheint richtig zu sein. Ihr Funktor erzeugt nur den Wert 9 für jedes Element, und Sie reduzieren 9 über 4 Elemente 9x4 = 36).

Nachdem all dies gesagt wurde, ist es mir nicht ganz klar, warum Schub den Hostpfad in Ihrem ursprünglichen Fall auswählt. Wenn Sie möchten, können Sie dafür eine,597.einreichen. Aber es ist durchaus möglich, dass ich nicht sorgfältig genug über das Schubabfertigungssystem nachgedacht habe. Die Verteilung des Hostcodealgorithmus (transform_reduce) könnte etwas verwirrend sein, weil es nicht offensichtlich ist, ob Sie beispielsweise Host- oder Gerätecontainer verwenden.