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 imbriquées, par exemple:

Dictionary<float, int> noOfNumbers = new Dictionary<float, int>();

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#, fonctionne sur un processeur Intel et nécessite plusieurs heures pour terminer. Je n'ai aucune connaissance des concepts de programmation GPU et des API, donc mes questions sont:

  • est-ce possible (et fait-il sens) d'utiliser un GPU pour accélérer de tels calculs?
  • Si oui: Est-ce que quelqu'un connaît un tutoriel ou a un exemple de code (le langage de programmation n'a pas d'importance)?

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

41
demandé sur dreamcrash 2012-11-09 06:33:44

5 réponses

Mise à jour 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;                      // e.g 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-ci, je viens de tester pour les entrées plus petites, parce que je teste mon ordinateur portable. Néanmoins, il a fait un travail. Cependant, il faut faire avancer les testicules.

Mise à jour Version séquentielle

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

Fondamentalement, il trie votre tableau de flotteurs. Il sera du voyage sur le tableau trié, analyser le nombre de fois qu'une valeur apparaît consécutivement dans le tableau, puis mettre cette valeur dans un dictionnaire avec le nombre de fois qu'il apparaît.

Vous pouvez utiliser la carte triée, au lieu de la unordered_map que j'ai utilisée.

Voici le code:

#include <stdio.h>
#include <stdlib.h>
#include "cuda.h"
#include <algorithm>
#include <string>
#include <iostream>
#include <tr1/unordered_map>


typedef std::tr1::unordered_map<float, int> 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<float, int> fill_dict(float *data, int size)
{
    float previous = data[0];
    int count = 1;
    std::tr1::unordered_map<float, int> 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<float, int> dict)
{
   for(std::tr1::unordered_map<float, int>::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<float, int> dict;

  generator(data,size);

  sort(data, data + size);
  dict = fill_dict(data,size);

  return 0;
}

Si vous avez la poussée de bibliothèque installée dans votre machine, vous devriez utiliser ceci:

#include <thrust/sort.h>
thrust::sort(data, data + size);

Au Lieu de cela,

sort(data, data + size);

C'est sûr que ce sera plus rapide.

D'Origine Post

"je travaille sur une application statistique qui a un grand tableau contenant 10 à 30 millions de valeurs à virgule flottante".

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

Oui, ça l'est. Il y a un mois, j'ai mis une simulation dynamique moléculaire entièrement sur le GPU. L'un des noyaux, qui calcule la force entre les paires de particules, recevoir 6 tableau chacun avec 500 000 doubles, un total de 3 millions doubles (22 MO).

Donc, vous prévoyez de mettre 30 Millions de points flottants c'est 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 dynamique moléculaire (MD), je dis non. La version séquentielle MD prend environ 25 heures pour terminer tandis que dans le GPU a pris 45 Minutes. Vous avez dit que votre application a pris quelques heures, également basé dans votre exemple de code, il semble plus doux que le Dynamique 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 d'un code dans Cuda pourrait être la somme de deux tableaux 2D:

En c:

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

Dans 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];
}

Dans Cuda, vous avez essentiellement pris chacun pour l'itération et divisé par chaque thread,

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

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

1) vous Donne le pour itération que chaque thread calculera en fonction de son id et de l'ID de bloc dans lequel se trouve le thread, le blockDim.x est le nombre de threads d'un bloc ont.

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 qui 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 parce que différents threads de différents blocs peuvent écrire / lire noOfOccurrences en même temps, donc vous devez incertain mutuel exclusion.

Ce n'est qu'une approche que vous pouvez même donner les itérations de la boucle externe aux threads au lieu des blocs.

Tutoriels

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

Et autres:

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

OpenCL: OpenCL Tutoriels / MacResearch

76
répondu dreamcrash 2012-11-17 22:19:00

Je ne connais rien au traitement parallèle ou au GPGPU, mais pour cet exemple spécifique, vous pourriez gagner beaucoup de temps en faisant un seul passage sur le tableau d'entrée plutôt que de le boucler un million de fois. Avec de grands ensembles de données, vous voudrez généralement faire les choses en un seul passage si possible. Même si vous faites plusieurs calculs indépendants, si c'est sur le même ensemble de données, vous pourriez obtenir une meilleure vitesse en les faisant tous dans la même passe, car vous obtiendrez une meilleure localité de référence de cette façon. Mais cela ne vaut peut-être pas la peine pour la complexité accrue de votre code.

En outre, vous ne voulez vraiment pas ajouter une petite quantité à un nombre à virgule flottante de manière répétitive comme ça, l'erreur d'arrondi s'additionnera 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 passage de votre échantillon le ferait regardez quelque chose comme ceci:

Dictionary<float, int> noOfNumbers = new Dictionary<float, int>();

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 aide.

11
répondu AlliedEnvy 2012-11-09 03:37:54

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

  • certainement Oui , ce type d'algorithme est généralement le candidat idéal pour le traitement parallélisme de données massif, les GPU sont si bons.

Si oui: Est-ce que quelqu'un connaît un tutoriel ou a obtenu un exemple de code (langage de programmation n'a pas d'importance)?

  • Lorsque vous voulez aller de la manière GPGPU, vous avez deux solutions de rechange: 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, et les processeurs aussi. Donc, vous devriez vraiment favorables.

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

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

  • Lorsque vous utilisez C# , vous pouvez utiliser des liaisons comme OpenCL.Net ou Cloo .

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

8
répondu Pragmateek 2012-11-13 10:20:17

En plus de la suggestion de l'affiche ci-dessus, utilisez le TPL (task parallel library) le cas échéant pour fonctionner en parallèle sur plusieurs cœurs.

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

Je ne sais pas si tous vos calculs correspondent correctement au GPU capacités, mais vous devrez utiliser un algorithme map-reduce de toute façon pour mapper les calculs sur les cœurs GPU, puis réduire les résultats partiels à un seul résultat, de sorte que vous pourriez aussi bien le faire sur le CPU avant de passer à une plate-forme moins familière.

6
répondu Eli Algranti 2012-11-09 03:49:05

Je ne suis pas sûr que l'utilisation de Gpu serait une bonne correspondance étant donné que les valeurs' largerFloatingPointArray ' doivent être récupérées de la mémoire. Ma compréhension est que les GPU sont mieux adaptés pour les calculs autonomes.

Je pense que transformer cette application de processus unique en une application distribuée fonctionnant sur de nombreux systèmes et peaufiner l'algorithme devrait accélérer considérablement les choses, en fonction du nombre de systèmes disponibles.

Vous pouvez utiliser le classique "diviser pour régner" approche. 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 base de données. 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. Le pire scénario est que chaque valeur ne se produit qu'une seule fois, mais c'est peu probable. Si largeFloatingPointArray continue de changer chaque fois que l'application est exécutée, la table de hachage en mémoire a du sens. Si elle est statique, la table peut être enregistrée dans une base de données clé-valeur telle que Berkeley DB. Appelons cela un système de "recherche".

Sur un autre système, appelons-le "main", créons des morceaux de travail et "dispersons" les éléments de travail sur N systèmes, et "rassemblons" les résultats à mesure qu'ils deviennent disponibles. Par exemple, un élément de travail pourrait être aussi simple que deux chiffres indiquant la plage sur laquelle un système devrait fonctionner. Lorsqu'un système termine le travail, il renvoie un tableau d'occurrences et il est prêt à travailler un autre morceau de travail.

La performance est améliorée car nous ne continuons pas à itérer sur largeFloatingPointArray. Si le système de recherche devient un goulot d'étranglement, il pourrait ê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 ciblé pour les systèmes basés sur plusieurs cœurs, souvent appelés microservers, qui sont/ou seront construits en utilisant plusieurs modules 'system-on-a-chip' dans un système. Les fournisseurs de modules ARM incluent Calxeda, AMD, AMCC, etc. Intel aura probablement également une offre similaire.

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

Si vous avez besoin d'aide pour implémenter une solution distribuée, je serais heureux d'en discuter avec vous.

Ajouté Le 16 Novembre 2012.

J'ai pensé un peu plus sur l'algorithme et je pense que cela devrait le faire en un seul passage. Il est écrit en C et il 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.
     *
     * E.g. 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]++;
        }   
    }
}
6
répondu Arun Taylor 2012-11-16 19:12:11