Puis-je / dois-je exécuter ce code sur un GPU?

Je travaille sur une application statistique contenant environ 10 à 30 millions de valeurs à virgule flottante dans un tableau.

Plusieurs méthodes effectuant des calculs différents mais indépendants sur le tableau dans des boucles nestedes, par exemple:

Dictionary noOfNumbers = new Dictionary(); for (float x = 0f; x < 100f; x += 0.0001f) { int noOfOccurrences = 0; foreach (float y in largeFloatingPointArray) { if (x == y) { noOfOccurrences++; } } noOfNumbers.Add(x, noOfOccurrences); } 

L’application en cours est écrite en C #, s’exécute sur un processeur Intel et nécessite plusieurs heures. Je n’ai aucune connaissance des concepts de programmation GPU et des API, mes questions sont donc les suivantes:

  • Est-il possible (et est-ce logique) d’utiliser un GPU pour accélérer ces calculs?
  • Si oui: Quelqu’un connaît-il un tutoriel ou a-t-il un exemple de code (le langage de programmation importe peu)?

Toute aide serait très appréciée.

UPDATE Version GPU

 __global__ void hash (float *largeFloatingPointArray,int largeFloatingPointArraySize, int *dictionary, int size, int num_blocks) { int x = (threadIdx.x + blockIdx.x * blockDim.x); // Each thread of each block will float y; // compute one (or more) floats int noOfOccurrences = 0; int a; while( x < size ) // While there is work to do each thread will: { dictionary[x] = 0; // Initialize the position in each it will work noOfOccurrences = 0; for(int j = 0 ;j < largeFloatingPointArraySize; j ++) // Search for floats { // that are equal // to it assign float y = largeFloatingPointArray[j]; // Take a candidate from the floats array y *= 10000; // eg if y = 0.0001f; a = y + 0.5; // a = 1 + 0.5 = 1; if (a == x) noOfOccurrences++; } dictionary[x] += noOfOccurrences; // Update in the dictionary // the number of times that the float appears x += blockDim.x * gridDim.x; // Update the position here the thread will work } } 

Celui que je viens de tester pour des entrées plus petites, parce que je teste mon ordinateur portable. Néanmoins, cela a fonctionné. Cependant, il est nécessaire de faire des testicules supplémentaires.

Version séquentielle UPDATE

Je viens de faire cette version naïve qui exécute votre algorithme pour 30 000 000 en moins de 20 secondes (compte déjà la fonction pour générer des données).

Fondamentalement, il sortinge votre tableau de flotteurs. Il parcourra le tableau sortingé, parsingra le nombre de fois qu'une valeur apparaîtra consécutivement dans le tableau, puis placera cette valeur dans un dictionnaire avec le nombre de fois où elle apparaît.

Vous pouvez utiliser une carte sortingée au lieu de la mappe unordered que j'ai utilisée.

Voici le code:

 #include  #include  #include "cuda.h" #include  #include  #include  #include  typedef std::tr1::unordered_map Mymap; void generator(float *data, long int size) { float LO = 0.0; float HI = 100.0; for(long int i = 0; i < size; i++) data[i] = LO + (float)rand()/((float)RAND_MAX/(HI-LO)); } void print_array(float *data, long int size) { for(long int i = 2; i < size; i++) printf("%f\n",data[i]); } std::tr1::unordered_map fill_dict(float *data, int size) { float previous = data[0]; int count = 1; std::tr1::unordered_map dict; for(long int i = 1; i < size; i++) { if(previous == data[i]) count++; else { dict.insert(Mymap::value_type(previous,count)); previous = data[i]; count = 1; } } dict.insert(Mymap::value_type(previous,count)); // add the last member return dict; } void printMAP(std::tr1::unordered_map dict) { for(std::tr1::unordered_map::iterator i = dict.begin(); i != dict.end(); i++) { std::cout << "key(string): " << i->first << ", value(int): " << i->second << std::endl; } } int main(int argc, char** argv) { int size = 1000000; if(argc > 1) size = atoi(argv[1]); printf("Size = %d",size); float data[size]; using namespace __gnu_cxx; std::tr1::unordered_map dict; generator(data,size); sort(data, data + size); dict = fill_dict(data,size); return 0; } 

Si vous avez installé la bibliothèque dans votre ordinateur, vous devez utiliser ceci:

 #include  thrust::sort(data, data + size); 

au lieu de cela

 sort(data, data + size); 

Bien sûr, ce sera plus rapide.

Poste d'origine

"Je travaille sur une application statistique qui contient un grand tableau contenant entre 10 et 30 millions de valeurs en virgule flottante".

"Est-il possible (et est-ce logique) d'utiliser un GPU pour accélérer ces calculs?"

Oui, ça l'est. Il y a un mois, je mets une simulation de dynamic moléculaire entièrement sur le GPU. L'un des kernelx, qui calcule la force entre les paires de particules, reçoit 6 tableaux chacun avec 500 000 doubles, soit un total de 3 millions de doubles (22 Mo).

Donc, vous prévoyez de mettre 30 millions de points flottants, ce qui représente environ 114 Mo de mémoire globale, donc ce n’est pas un problème, même mon ordinateur portable a 250 Mo.

Le nombre de calcul peut être un problème dans votre cas? Sur la base de mon expérience avec la dynamic moléculaire (MD), je dis non. La version séquentielle de MD prend environ 25 heures, tandis que celle en GPU prend 45 minutes. Vous avez dit que votre application prenait quelques heures, également dans votre exemple de code, elle semble plus douce que la dynamic moléculaire.

Voici l'exemple de calcul de force:

 __global__ void add(double *fx, double *fy, double *fz, double *x, double *y, double *z,...){ int pos = (threadIdx.x + blockIdx.x * blockDim.x); ... while(pos < particles) { for (i = 0; i < particles; i++) { if(//inside of the same radius) { // calculate force } } pos += blockDim.x * gridDim.x; } } 

Un exemple simple de code dans Cuda pourrait être la sum de deux tableaux 2D:

En c:

 for(int i = 0; i < N; i++) c[i] = a[i] + b[i]; 

À Cuda:

 __global__ add(int *c, int *a, int*b, int N) { int pos = (threadIdx.x + blockIdx.x) for(; i < N; pos +=blockDim.x) c[pos] = a[pos] + b[pos]; } 

En Cuda, vous avez pris chacun pour itération et diviser par chaque thread,

 1) threadIdx.x + blockIdx.x*blockDim.x; 

Chaque bloc a un identifiant de 0 à N-1 (N le nombre maximum de blocs) et chaque bloc a un nombre X de threads avec un identifiant de 0 à X-1.

1) Vous donne pour itération que chaque thread calculera en fonction de son identifiant et de l'identifiant du bloc dans lequel se trouve le thread, le blockDim.x est le nombre de threads d'un bloc.

Donc, si vous avez 2 blocs chacun avec 10 threads et un N = 40, le:

 Thread 0 Block 0 will execute pos 0 Thread 1 Block 0 will execute pos 1 ... Thread 9 Block 0 will execute pos 9 Thread 0 Block 1 will execute pos 10 .... Thread 9 Block 1 will execute pos 19 Thread 0 Block 0 will execute pos 20 ... Thread 0 Block 1 will execute pos 30 Thread 9 Block 1 will execute pos 39 

En regardant votre code, j'ai fait ce brouillon de ce que cela pourrait être dans cuda:

 __global__ hash (float *largeFloatingPointArray, int *dictionary) // You can turn the dictionary in one array of int // here each position will represent the float // Since x = 0f; x < 100f; x += 0.0001f // you can associate each x to different position // in the dictionary: // pos 0 have the same meaning as 0f; // pos 1 means float 0.0001f // pos 2 means float 0.0002f ect. // Then you use the int of each position // to count how many times that "float" had appeared int x = blockIdx.x; // Each block will take a different x to work float y; while( x < 1000000) // x < 100f (for incremental step of 0.0001f) { int noOfOccurrences = 0; float z = converting_int_to_float(x); // This function will convert the x to the // float like you use (x / 0.0001) // each thread of each block // will takes the y from the array of largeFloatingPointArray for(j = threadIdx.x; j < largeFloatingPointArraySize; j += blockDim.x) { y = largeFloatingPointArray[j]; if (z == y) { noOfOccurrences++; } } if(threadIdx.x == 0) // Thread master will update the values atomicAdd(&dictionary[x], noOfOccurrences); __syncthreads(); } 

Vous devez utiliser atomicAdd car différents threads provenant de différents blocs peuvent écrire / lire des noOfOccurrences en même temps, vous devez donc vous assurer de l'exclusion mutuelle.

Ceci est seulement une approche possible, vous pouvez même donner les itérations de la boucle externe aux threads au lieu des blocs.

Des tutoriels

La série Dr Dobbs Journal CUDA: Supercomputing pour les masses par Rob Farmer est excellente et couvre à peu près tout dans ses quatorze versements. Il commence aussi plutôt doucement et est donc assez amical pour les débutants.

et d'autres:

  • Développer avec CUDA - Introduction
  • Volume I: Introduction à la programmation CUDA
  • Démarrer avec CUDA
  • Liste des ressources CUDA

Jetez un oeil sur le dernier élément, vous trouverez de nombreux liens pour apprendre CUDA.

OpenCL: Didacticiels OpenCL | MacResearch

Je ne connais pas grand chose au parallel processing ou au GPGPU, mais pour cet exemple spécifique, vous pourriez gagner beaucoup de temps en effectuant un seul passage sur le tableau d’entrée plutôt que de le parcourir un million de fois. Avec des ensembles de données volumineux, vous voudrez généralement faire les choses en un seul passage si possible. Même si vous effectuez plusieurs calculs indépendants, si vous effectuez le même dataset, vous obtiendrez une meilleure vitesse tout en ayant la même localisation de référence. Mais cela ne vaut peut-être pas la peine pour la complexité accrue de votre code.

En outre, vous ne voulez vraiment pas append une petite quantité à un nombre à virgule flottante de manière répétitive, l’erreur d’arrondi s’ajoutant et vous n’obtiendrez pas ce que vous vouliez. J’ai ajouté une instruction if à mon exemple ci-dessous pour vérifier si les entrées correspondent à votre modèle d’itération, mais omettez-le si vous n’en avez pas réellement besoin.

Je ne connais aucun C #, mais une seule implémentation de votre échantillon ressemblerait à ceci:

 Dictionary noOfNumbers = new Dictionary(); foreach (float x in largeFloatingPointArray) { if (math.Truncate(x/0.0001f)*0.0001f == x) { if (noOfNumbers.ContainsKey(x)) noOfNumbers.Add(x, noOfNumbers[x]+1); else noOfNumbers.Add(x, 1); } } 

J’espère que cela t’aides.

Est-il possible (et est-ce logique) d’utiliser un GPU pour accélérer ces calculs?

  • Certainement OUI , ce type d’algorithme est généralement le candidat idéal pour le traitement massif du parallélisme de données, ce à quoi les GPU sont si efficaces.

Si oui: Quelqu’un connaît-il un tutoriel ou a-t-il un exemple de code (le langage de programmation importe peu)?

  • Lorsque vous voulez utiliser la méthode GPGPU, vous avez deux alternatives: CUDA ou OpenCL .

    CUDA est mature avec beaucoup d’outils mais est centré sur les GPU NVidia.

    OpenCL est un standard fonctionnant sur les GPU NVidia et AMD, ainsi que sur les processeurs. Donc, vous devriez vraiment le favoriser.

  • Pour le tutoriel, vous avez une excellente série sur CodeProject de Rob Farber : http://www.codeproject.com/Articles/Rob-Farber#Articles

  • Pour votre cas d’utilisation spécifique, il y a beaucoup d’échantillons pour les histogrammes construits avec OpenCL (notez que beaucoup sont des histogrammes mais les principes sont les mêmes).

  • En utilisant C #, vous pouvez utiliser des liaisons telles que OpenCL.Net ou Cloo .

  • Si votre tableau est trop gros pour être stocké dans la mémoire GPU, vous pouvez le partitionner et réexécuter facilement votre kernel OpenCL pour chaque partie.

En plus de la suggestion de l’affiche ci-dessus, utilisez la bibliothèque parallèle de tâches (TPL) lorsque cela est approprié pour s’exécuter en parallèle sur plusieurs cœurs.

L’exemple ci-dessus pourrait utiliser Parallel.Foreach et ConcurrentDictionary, mais une configuration de réduction de carte plus complexe où le tableau est divisé en morceaux générant chacun un dictionnaire qui serait ensuite réduit à un dictionnaire unique vous donnerait de meilleurs résultats.

Je ne sais pas si tous vos calculs correspondent correctement aux capacités du GPU, mais vous devrez utiliser un algorithme de réduction de carte pour mapper les calculs aux cœurs du GPU, puis réduire les résultats partiels à un seul résultat. autant le faire sur le processeur avant de passer à une plate-forme moins familière.

Je ne suis pas sûr que l’utilisation des GPU soit un bon résultat, étant donné que les valeurs “largerFloatingPointArray” doivent être récupérées de la mémoire. D’après ce que je comprends, les GPU conviennent mieux aux calculs autonomes.

Je pense que transformer cette application de processus unique en une application dissortingbuée exécutée sur de nombreux systèmes et modifier l’algorithme devrait accélérer considérablement les choses, en fonction du nombre de systèmes disponibles.

Vous pouvez utiliser l’approche classique de «diviser pour régner». L’approche générale que je prendrais est la suivante.

Utilisez un système pour prétraiter «largeFloatingPointArray» dans une table de hachage ou une firebase database. Cela se ferait en un seul passage. Il utiliserait la valeur à virgule flottante comme clé et le nombre d’occurrences dans le tableau comme valeur. Dans le pire des cas, chaque valeur ne se produit qu’une fois, mais cela est peu probable. Si grandFloatingPointArray continue à changer chaque fois que l’application est exécutée, la table de hachage en mémoire est logique. S’il est statique, la table peut être enregistrée dans une firebase database de valeurs-clés telle que Berkeley DB. Appelons cela un système de recherche.

Sur un autre système, appelons-le «principal», créez des morceaux de travail et «dispersez» les éléments de travail dans N systèmes, et «collectez» les résultats dès qu’ils sont disponibles. Par exemple, un élément de travail peut être aussi simple que deux chiffres indiquant la plage sur laquelle un système doit fonctionner. Lorsqu’un système termine le travail, il renvoie des occurrences et il est prêt à travailler sur une autre partie du travail.

Les performances sont améliorées car nous ne continuons pas à itérer sur largeFloatingPointArray. Si le système de consultation devient un goulot d’étranglement, il peut être répliqué sur autant de systèmes que nécessaire.

Avec un nombre suffisant de systèmes fonctionnant en parallèle, il devrait être possible de réduire le temps de traitement à quelques minutes.

Je travaille sur un compilateur pour la programmation parallèle en C ciblant les systèmes basés sur plusieurs cœurs, souvent appelés microserveurs, qui sont ou seront construits en utilisant plusieurs modules «système sur puce» au sein d’un système. Les fournisseurs de modules ARM incluent Calxeda, AMD, AMCC, etc. Intel offrira probablement une offre similaire.

J’ai une version du compilateur qui peut être utilisée pour une telle application. Le compilateur, basé sur des prototypes de fonctions C, génère un code de mise en réseau C qui implémente le code de communication inter-processus (IPC) entre les systèmes. L’un des mécanismes IPC disponibles est socket / tcp / ip.

Si vous avez besoin d’aide pour mettre en place une solution dissortingbuée, je serais ravi d’en discuter avec vous.

Ajouté le 16 novembre 2012.

J’ai réfléchi un peu plus à l’algorithme et je pense que cela devrait se faire en un seul passage. Il est écrit en C et devrait être très rapide par rapport à ce que vous avez.

 /* * Convert the X range from 0f to 100f in steps of 0.0001f * into a range of integers 0 to 1 + (100 * 10000) to use as an * index into an array. */ #define X_MAX (1 + (100 * 10000)) /* * Number of floats in largeFloatingPointArray needs to be defined * below to be whatever your value is. */ #define LARGE_ARRAY_MAX (1000) main() { int j, y, *noOfOccurances; float *largeFloatingPointArray; /* * Allocate memory for largeFloatingPointArray and populate it. */ largeFloatingPointArray = (float *)malloc(LARGE_ARRAY_MAX * sizeof(float)); if (largeFloatingPointArray == 0) { printf("out of memory\n"); exit(1); } /* * Allocate memory to hold noOfOccurances. The index/10000 is the * the floating point number. The contents is the count. * * Eg noOfOccurances[12345] = 20, means 1.2345f occurs 20 times * in largeFloatingPointArray. */ noOfOccurances = (int *)calloc(X_MAX, sizeof(int)); if (noOfOccurances == 0) { printf("out of memory\n"); exit(1); } for (j = 0; j < LARGE_ARRAY_MAX; j++) { y = (int)(largeFloatingPointArray[j] * 10000); if (y >= 0 && y <= X_MAX) { noOfOccurances[y]++; } } }