CUDA Ajouter des lignes d'une matrice

Question

J'essaie d'ajouter les lignes d'une matrice 4800x9600 ensemble, ce qui entraîne une matrice 1x9600.

Ce que j'ai fait est divisé le 4800x9600 en 9.600 matrices de longueur 4800 chacun. J'effectue ensuite une réduction sur les 4800 éléments.

Le problème est, ce qui est vraiment lent ...

Quelqu'un at-il des suggestions? Fondamentalement, j'essaie d'implémenter la fonction somme (...) de MATLAB.

Voici le code que j'ai vérifié fonctionne très bien, il est juste qu'il est vraiment lent:

void reduceRows(Matrix Dresult,Matrix DA) 
{ 
     //split DA into chunks 
     Matrix Dchunk; 
     Dchunk.h=1;Dchunk.w=DA.h; 
     cudaMalloc((void**)&Dchunk.data,Dchunk.h*Dchunk.w*sizeof(float)); 

     Matrix DcolSum; 
     DcolSum.h=1;DcolSum.w=1; 
     //cudaMalloc((void**)&DcolSum.data,DcolSum.h*DcolSum.w*sizeof(float)); 

     int i; 
     for(i=0;i<DA.w;i++) //loop over each column 
     { 
       //printf("%d ",i); 
       cudaMemcpy(Dchunk.data,&DA.data[i*DA.h],DA.h*sizeof(float),cudaMemcpyDeviceToDevice); 
       DcolSum.data=&Dresult.data[i]; 
       reduceTotal(DcolSum,Dchunk); 
     } 
     cudaFree(Dchunk.data); 
} 

Matrice est définie comme:

typedef struct{ 
     long w; 
     long h; 
     float* data; 
}Matrix; 

ReduceTotal() appelle simplement la norme réduction NVIDIA , résume tous les éléments dans Dchunk et met la réponse dans DcolSum.

Je suis sur le point de faire tout cela sur la CPU si je ne peux pas trouver une réponse ...; (

Un grand merci à l'avance,

Answer 1

Au lieu de boucler sur chaque colonne, parallélisez sur les colonnes. Chacun des 4600 threads résume les 9600 entrées dans sa colonne, et met la somme à l'endroit approprié dans le vecteur de résultat.

Si vous êtes à la recherche d'une bibliothèque pour faire travailler avec Cuda plus simple, je recommande fortement POUSSÉE: http://code.google.com/p/thrust/

utilisant la poussée, je créerais un foncteur de tenir votre pointeur de matrice dans la mémoire de l'appareil et la carte puis sur une séquence d'indices de colonne. L'opérateur() du foncteur prendrait un index, résumerait tout dans cette colonne de la matrice, et renverrait la somme.Ensuite, vous auriez votre somme assis dans un thrust :: device_vector sans aucune copie de la mémoire (ou même des appels CUDA directs).

Votre foncteur pourrait ressembler à:

struct ColumnSumFunctor { 
    const Matrix matrix; 

    // Make a functor to sum the matrix 
    InitialGuessFunctor(const Matrix& matrix); 

    // Compute and return the sum of the specified column 
    __device__ 
    int operator()(const int& column) const; 
}; 

Answer 2

réduction est très basique dans le fonctionnement GPGPU, il est censé être rapide, et 9600 fois la réduction ne devrait pas être lent non plus.

Quelle carte graphique utilisez-vous

? Je vous suggère de le diviser en 9600 tableaux, chaque fois que vous réduisez un tableau de 4800 éléments en un résultat, au lieu de réduireTotal, je vous suggère se CUDPP pour effectuer l'opération de réduction, CUDPP est comme le STL pour CUDA. Il est mis en œuvre avec inquiétude sur la performance.

http://code.google.com/p/cudpp/

Answer 3

Je pense que votre problème est que vous lancez des noyaux de 9600X2. Cela devrait être un algorithme facile à exprimer en tant que noyau unique.

La manière la plus naïve de l'implémenter ne fusionnerait pas la mémoire, mais elle pourrait bien être plus rapide que la façon dont vous le faites maintenant.

Une fois que vous avez la façon naïve de travailler, alors fusionnez votre mémoire: par ex. chaque thread dans un bloc lit 16 flottants consécutifs dans la mémoire partagée, syncthreads, puis accumule les 16 flottants pertinents dans un registre, synthétise, puis répète

Le SDK Computing contient de nombreux exemples de techniques de réduction.