Wie erzwinge ich einen Funktor, um einen ganzen Schub zu sehen :: Vektor, so dass Sortierung möglich ist?

Question

Ich bin neu in CUDA und habe ein paar Probleme mit Funktoren. Ich versuche einen Schub :: Vektor von Schub :: Vektoren in einen Funktor einzugeben. Zur Zeit kann ich einen Vektor eingeben und jedem Element etwas zufügen und den modifizierten Vektor mit Hilfe von schub :: for_each zurückgeben, aber wenn ich einen Vektor in einem Funktor sortieren möchte, müsste ich den gesamten Vektor gleichzeitig eingeben können Der Funktor kann als Ganzes darauf reagieren. Gibt es eine Möglichkeit, dies zu tun?

Der folgende Code kompiliert, aber gibt den Vektor nicht sortiert zurück.

#include <iostream> 
#include <fstream> 
#include <sstream> 
#include <string> 
#include <vector> 
#include <iterator> 
#include <stdlib.h> 
#include <cuda.h> 
#include <cuda_runtime.h> 
#include <device_launch_parameters.h> 
#include <thrust/functional.h> 
#include <thrust/device_vector.h> 
#include <thrust/host_vector.h> 
#include <thrust/reduce.h> 
#include <thrust/transform_reduce.h> 
#include <thrust/transform.h> 
#include <thrust/sort.h> 
#include <thrust/execution_policy.h> 
#include <thrust/system/cuda/execution_policy.h> 
#include <thrust/tuple.h> 
#include <thrust/count.h> 
#include <thrust/sequence.h> 
#include <thrust/iterator/zip_iterator.h> 
#include <thrust/for_each.h> 
#include <ctime> 
#include <cstdio> 
#include <cassert> 

using namespace std; 
template<typename T> 
struct sort_vector 
{ 
    __host__ __device__ thrust::device_vector<float> operator() (thrust::tuple<thrust::device_vector<float>, thrust::device_vector<float>>  x) 
    { 
    thrust::device_vector<float> y = thrust::get<0>(x); 
    thrust::sort(y.begin(), y.end()); 
    return thrust::get<1>(x) = y; 
    } 
}; 

int main() { 
    thrust::device_vector<float> d_fraction(5); 
    d_fraction[0] = 1; 
    d_fraction[1] = 5; 
    d_fraction[2] = 3; 
    d_fraction[3] = 2; 
    d_fraction[4] = 4; 

    cout << "original" << endl; 
    int f = 0; 
    while (f < 5){ 
     cout << d_fraction[f] << endl; 
     f++; 
    } 

    cudaStream_t s1; 
    cudaStreamCreate(&s1); 
    thrust::device_vector<float> result1(5); 

    thrust::for_each(thrust::cuda::par.on(s1), 
    thrust::make_zip_iterator(thrust::make_tuple(d_fraction.begin(), result1.begin())), 
    thrust::make_zip_iterator(thrust::make_tuple(d_fraction.end(), result1.end())), sort_vector<thrust::device_vector<float>>()); 

    cudaStreamSynchronize(s1); 

    cout << "sorted" << endl; 
    int d = 0; 
    while (d < 5){ 
     cout << Sresult2[d] << endl; 
     d++; 
    } 

    cudaStreamDestroy(s1); 
    return 0; 
} 

Allerdings, wenn ich versuche, eine Referenz zu verwenden, wie

_host__ __device__ thrust::device_vector<float> operator() (thrust::tuple<thrust::device_vector<float> &, thrust::device_vector<float> &>  x) 

Der Code nicht mehr kompiliert.

Ist es möglich, muss ich einen Referenzzeiger für den Vektor, so dass der Funktor den gesamten Vektor sehen kann? Oder ist es möglich, dass das Problem ist, dass ich den Vektor nach Wert übergebe und es gibt eine andere Art, die ich nicht kenne, einen Vektor in einen Funktor zu übergeben?

Answer 1

Ein Funktor arbeitet normalerweise aus dem Kontext eines einzelnen Threads. Wenn wir das CUDA-Backend verwenden, sprechen wir von einem CUDA-Thread.

Der typische Ansatz zum Sortieren eines Vektors wäre thrust::sort direkt auf dem Vektor zu verwenden. In den trivialsten Verwendungen wird überhaupt keine Funktordefinition benötigt.

Wenn Sie einen Vektor "innerhalb eines Funktors" sortieren wollen, dann ist es notwendig, einen Zeiger auf diesen Vektor zum Funktor zu übergeben und den Funktor daran arbeiten zu lassen.

Schubgerätecode (was im Kontext eines Funktors ausgeführt wird) kann im Allgemeinen nicht direkt mit Konstrukten wie thrust::device_vector umgehen. Es ist derzeit auch nicht möglich, einen Einrichtungsvektor von Einrichtungsvektoren zu konstruieren.

Deshalb habe ich Ihren Code zu etwas geändert, das praktikabel ist und "innerhalb eines Funktors" sortiert. Ich habe gewählt, die Vektoren zu verknüpfen, um in einen einzelnen Vektor sortiert zu werden. Wir geben die Adresse dieses Vektors an der Sortier Funktors und dann jedem Thread berechnet seine Reichweite zu sortieren, und übergibt die für sequentielle Sortierung im Thread zu thrust::sort:

$ cat t1211.cu 
#include <iostream> 
#include <thrust/device_vector.h> 
#include <thrust/host_vector.h> 
#include <thrust/sort.h> 
#include <thrust/execution_policy.h> 
#include <thrust/for_each.h> 
#include <thrust/sequence.h> 
#include <cstdlib> 

const int num_segs = 3; // number of segments to sort 
const int num_vals = 5; // number of values in each segment 
const int range = 100; // range of values 

using namespace std; 

template <typename T> 
struct sort_vector 
{ 
    T *data; 
    sort_vector(T *_data) : data(_data) {}; 
    __host__ __device__ void operator()(int idx) 
    { 
    thrust::sort(thrust::seq, data+idx*num_vals, data+((idx+1)*num_vals)); 
    } 
}; 

int main() { 
    thrust::device_vector<float> d_data(num_segs*num_vals); 
    for (int i = 0; i < num_segs*num_vals; i++) 
     d_data[i] = rand()%range; 

    cout << "original" << endl; 
    int f = 0; 
    while (f < num_segs*num_vals){ 
     cout << d_data[f] << endl; 
     f++; 
    } 
    thrust::device_vector<int> d_idxs(num_segs); 
    thrust::sequence(d_idxs.begin(), d_idxs.end()); 
    cudaStream_t s1; 
    cudaStreamCreate(&s1); 
    //thrust::device_vector<float> result1(5); 

    thrust::for_each(thrust::cuda::par.on(s1), 
    d_idxs.begin(), 
    d_idxs.end(), sort_vector<float>(thrust::raw_pointer_cast(d_data.data()))); 

    cudaStreamSynchronize(s1); 

    cout << "sorted" << endl; 
    int d = 0; 
    while (d < num_segs*num_vals){ 
     cout << d_data[d] << endl; 
     d++; 
    } 

    cudaStreamDestroy(s1); 
    return 0; 
} 
$ nvcc -o t1211 t1211.cu 
$ ./t1211 
original 
83 
86 
77 
15 
93 
35 
86 
92 
49 
21 
62 
27 
90 
59 
63 
sorted 
15 
77 
83 
86 
93 
21 
35 
49 
86 
92 
27 
59 
62 
63 
90 
$ 

In diesem Fall, wie ersichtlich ist durch thrust::seq , die Arbeit in jedem Thread wird sequentiell erledigt. (Zusammen arbeiten die Threads hier parallel, aber sie kooperieren nicht - jeder Thread arbeitet an einem unabhängigen Problem).

Dies ist nicht die einzige mögliche Lösung. Sie könnten interessiert sein an this Frage/Antwort, die eine Vielzahl von anderen verwandten Ideen hat.

Um klar zu sein, was ich denke, Sie diskutieren hier ist eine "vektorisierte" (oder segmentierte) Art. Dies ist nicht der schnellste Ansatz, aber ich versuche, einige praktikable Konzepte für Sie als eine direkte Erweiterung zu dem, was Sie gezeigt haben, zu demonstrieren, um Ihre Frage zu beantworten ("Wie man einen Funktor dazu bringt, einen ganzen Schubvektor zu sehen so dass das Sortieren möglich ist? ") Ein schnellerer Ansatz für die vektorisierte Sortierung wird in der oben verlinkten Frage/Antwort diskutiert.