Comment forcer un foncteur à voir un ensemble thrust :: vector pour que le tri soit possible?

Question

Je suis nouveau à CUDA et j'ai un peu de mal avec les foncteurs. J'essaie d'entrer un thrust :: vector of thrust :: vectors dans un foncteur. Actuellement, je peux entrer un vecteur et faire quelque chose à chaque élément et retourner le vecteur modifié en utilisant thrust :: for_each, mais si je devais trier un vecteur dans un foncteur, je devrais être capable d'entrer le vecteur entier le foncteur peut agir sur lui dans son ensemble. Y a-t-il un moyen de faire cela?

Le code ci-dessous compile, mais ne renvoie pas le vecteur trié.

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

Cependant, lorsque je tente d'utiliser une référence telle que

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

Le code ne compile.

Est-il possible que je doive lancer un pointeur de référence pour le vecteur afin que le foncteur puisse voir le vecteur entier? Ou est-ce possible que le problème est que je passe le vecteur par la valeur et il y a une façon différente que je ne connais pas de passer un vecteur dans un foncteur?

Answer 1

Un foncteur fonctionne normalement à partir du contexte d'un seul thread. Si vous utilisez le backend CUDA, nous parlons d'un thread CUDA.

L'approche typique pour trier un vecteur serait d'utiliser thrust::sort directement sur le vecteur. Dans les usages les plus triviaux, aucune définition de foncteur n'est nécessaire.

Si vous vouliez trier un vecteur "à l'intérieur d'un foncteur", il serait alors nécessaire de passer un pointeur vers ce vecteur au foncteur, et laisser le foncteur travailler dessus.

Le code de périphérique de poussée (ce qui est exécuté dans le contexte d'un foncteur) ne peut généralement pas gérer les constructions comme thrust::device_vector directement. Il n'est pas non plus possible actuellement de construire un vecteur de périphérique de vecteurs de périphérique. Par conséquent j'ai modifié votre code pour quelque chose qui est réalisable et trie "dans un functor". J'ai choisi de concaténer les vecteurs à trier en un seul vecteur. Nous passons l'adresse de ce vecteur au foncteur de tri, puis chaque fil calcule sa gamme pour trier et passes à thrust::sort pour le tri séquentiel dans le fil:

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

Dans ce cas, comme on le voit par thrust::seq , le travail effectué dans chaque thread se fait séquentiellement. (Ensemble, les threads fonctionnent ici en parallèle, mais ils ne coopèrent pas - chaque thread travaille sur un problème indépendant).

Ce n'est pas la seule solution possible. Vous pouvez être intéressé par this question/réponse qui a une variété d'autres idées connexes. Pour être clair, je pense que vous parlez ici d'un tri «vectorisé» (ou segmenté). Ce n'est pas l'approche la plus rapide, mais j'essaie de démontrer quelques concepts réalisables pour vous comme une simple extension à ce que vous avez montré, afin de répondre à votre question ("Comment forcer un foncteur à voir un vecteur entier ») Une approche plus rapide pour le tri vectorisé est discutée dans la question/réponse liée ci-dessus.