c++ - tutorial - Cuda kernel regresando vectores

parallel programming cuda

Question

Tengo una lista de palabras, mi objetivo es hacer coincidir cada palabra en una frase muy larga. No tengo problemas para emparejar cada palabra, mi único problema es devolver un vector de estructuras que contenga información sobre cada coincidencia.

En codigo:

typedef struct {
int A, B, C; } Match;
__global__ void Find(veryLongPhrase * _phrase, Words * _word_list, vector<Match> * _matches)
{
int a, b, c;
[...] //Parallel search for each word in the phrase
if(match) //When an occurrence is found
{
_matches.push_back(new Match{ A = a, B = b, C = c }); //Here comes the unknown, what should I do here???
}
}
main()
{
[...]
veryLongPhrase * myPhrase = "The quick brown fox jumps over the lazy dog etc etc etc..."
Words * wordList = {"the", "lazy"};
vector<Match> * matches; //Obviously I can''t pass a vector to a kernel
Find<<< X, Y >>>(myPhrase, wordList, matches);
[...]
}

He probado la biblioteca de Thrust pero sin éxito, ¿me pueden sugerir algún tipo de solución?

Muchas gracias.

Answer 1

algo como esto debería funcionar (codificado en el navegador, no probado):

// N is the maximum number of structs to insert
#define N 10000
typedef struct {
int A, B, C; } Match;
__device__ Match dev_data[N];
__device__ int dev_count = 0;
__device__ int my_push_back(Match * mt) {
int insert_pt = atomicAdd(&dev_count, 1);
if (insert_pt < N){
dev_data[insert_pt] = *mt;
return insert_pt;}
else return -1;}
__global__ void Find(veryLongPhrase * _phrase, Words * _word_list, vector<Match> * _matches)
{
int a, b, c;
[...] //Parallel search for each word in the phrase
if(match) //When an occurrence is found
{
my_push_back(new Match{ A = a, B = b, C = c }); }
}
main()
{
[...]
veryLongPhrase * myPhrase = "The quick brown fox jumps over the lazy dog etc etc etc..."
Words * wordList = {"the", "lazy"};
Find<<< X, Y >>>(myPhrase, wordList);
int dsize;
cudaMemcpyFromSymbol(&dsize, dev_count, sizeof(int));
vector<Match> results(dsize);
cudaMemcpyFromSymbol(&(results[0]), dev_data, dsize*sizeof(Match));
[...]
}

Esto requerirá capacidad de cálculo 1.1 o superior para la operación atómica.

nvcc -arch=sm_11 ...

Aquí hay un ejemplo trabajado:

$ cat t347.cu
#include <iostream>
#include <vector>
// N is the maximum number of structs to insert
#define N 10000
typedef struct {
int A, B, C; } Match;
__device__ Match dev_data[N];
__device__ int dev_count = 0;
__device__ int my_push_back(Match & mt) {
int insert_pt = atomicAdd(&dev_count, 1);
if (insert_pt < N){
dev_data[insert_pt] = mt;
return insert_pt;}
else return -1;}
__global__ void Find()
{
if(threadIdx.x < 10) //Simulate a found occurrence
{
Match a = { .A = 1, .B = 2, .C = 3 };
my_push_back(a); }
}
main()
{
Find<<< 2, 256 >>>();
int dsize;
cudaMemcpyFromSymbol(&dsize, dev_count, sizeof(int));
if (dsize >= N) {printf("overflow error/n"); return 1;}
std::vector<Match> results(dsize);
cudaMemcpyFromSymbol(&(results[0]), dev_data, dsize*sizeof(Match));
std::cout << "number of matches = " << dsize << std::endl;
std::cout << "A = " << results[dsize-1].A << std:: endl;
std::cout << "B = " << results[dsize-1].B << std:: endl;
std::cout << "C = " << results[dsize-1].C << std:: endl;
}
$ nvcc -arch=sm_11 -o t347 t347.cu
$ ./t347
number of matches = 20
A = 1
B = 2
C = 3
$

Tenga en cuenta que en este caso mi creación de estructura de resultado de Match es diferente, y estoy pasando por referencia, pero el concepto es el mismo.