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

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:

• Développer Avec CUDA-Introduction
• Volume I: Introduction à la programmation CUDA
• mise en route avec CUDA
• liste des ressources CUDA

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

OpenCL: OpenCL Tutoriels / MacResearch